wangdequan/workinf_Blender_Wasm
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Complete V1 performance and OOM release gates

Browse Source
This commit is contained in:
mes123456
2026-08-14 22:32:09 -04:00
parent 3ea9974eee
commit a3f3071c03
45 changed files with 4206 additions and 276 deletions
Download Patch File Download Diff File Expand all files Collapse all files

259
web/app/src/render/nanovdb-volume-renderer.ts
View File

@@ -5,6 +5,7 @@ export interface NanoVDBWebGPUCapabilityIR {
reason?: string;
maxStorageBufferBindingSize?: number;
maxBufferSize?: number;
maxStorageBuffersPerShaderStage?: number;
}
export interface NanoVDBWebGPUGrid {
@@ -15,14 +16,20 @@ export interface NanoVDBWebGPUGrid {
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;
evictPage(pageIndex: number): void;
hasResidentPage(pageIndex: number): boolean;
dispose(): void;
}
export type NanoVDBViewAxis = "X" | "Y" | "Z";
export interface NanoVDBMaterialGridUploadsIR {
temperature?: NanoVDBWebGPUGrid;
color?: NanoVDBWebGPUGrid;
@@ -156,18 +163,20 @@ function specializeFloatTraversal(prefix: string, gridName: string, pageTableNam
return source;
}
const temperatureTraversalWGSL = specializeFloatTraversal("temperature", "temperature_grid", "", "temperature");
const emissionTraversalWGSL = specializeFloatTraversal("emission", "emission_grid", "", "emission");
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) || params.color_page_bytes == 0u) { return 0u; }
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 = page;
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; }
@@ -227,15 +236,18 @@ fn sample_color_linear(position: vec3<f32>) -> vec4<f32> {
}
`;
export async function probeNanoVDBWebGPU(requiredBytes = 1): Promise<{ capability: NanoVDBWebGPUCapabilityIR; adapter?: GPUAdapter; device?: GPUDevice }> {
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);
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 };
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 {
@@ -254,11 +266,15 @@ export function uploadNanoVDBFloat32Grid(device: GPUDevice, payload: ArrayBuffer
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,
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(); },
@@ -271,61 +287,108 @@ export function uploadNanoVDBFloat32GridPaged(
pageByteLength: number,
maxResidentBytes: number,
): NanoVDBWebGPUGrid {
if (payload.byteLength === 0 || payload.byteLength % 32 !== 0 || !Number.isSafeInteger(pageByteLength) || pageByteLength < 64 * 1024 || pageByteLength % 32 !== 0) {
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(payload.byteLength / pageByteLength);
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");
}
const buffer = device.createBuffer({ label: "NanoVDB paged Float32 grid", size: physicalBytes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
let buffer: GPUBuffer | undefined;
let pageTable: GPUBuffer | undefined;
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();
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 });
new Uint32Array(pageTable.getMappedRange()).fill(0xffffffff);
pageTable.unmap();
}
catch (error) {
pageTable?.destroy();
buffer?.destroy();
throw error;
}
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;
const lastUsed = new Map<number, number>();
let clock = 0;
let evictions = 0;
const writePageTable = (pageIndex: number, slot: number): void => {
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);
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].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");
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: payload.byteLength,
byteLength,
pageByteLength,
pageCount,
residentPageCount: resident.size,
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,
evictPage: (pageIndex) => {
if (!resident.delete(pageIndex)) return;
table[pageIndex] = 0xffffffff;
device.queue.writeBuffer(pageTable, pageIndex * 4, new Uint32Array([0xffffffff]));
touchPage: (pageIndex) => {
if (!resident.has(pageIndex)) return false;
lastUsed.set(pageIndex, ++clock);
return true;
},
evictPage: evict,
hasResidentPage: (pageIndex) => resident.has(pageIndex),
dispose: () => { resident.clear(); buffer.destroy(); pageTable.destroy(); },
dispose: () => { resident.clear(); lastUsed.clear(); buffer.destroy(); pageTable.destroy(); },
};
}
@@ -398,9 +461,9 @@ export class NanoVDBWebGPUDeviceSession {
return () => this.lossListeners.delete(listener);
}
async open(requiredBytes: number): Promise<GPUDevice> {
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);
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++;
@@ -418,9 +481,9 @@ export class NanoVDBWebGPUDeviceSession {
return this.loss;
}
async recover(requiredBytes: number): Promise<GPUDevice> {
async recover(requiredBytes: number, requiredStorageBuffers = 5): Promise<GPUDevice> {
this.device?.destroy();
return this.open(requiredBytes);
return this.open(requiredBytes, requiredStorageBuffers);
}
dispose(): void {
@@ -436,6 +499,47 @@ function paramsBuffer(device: GPUDevice, values: Uint32Array): GPUBuffer {
return buffer;
}
export async function readNanoVDBWordsWebGPU(device: GPUDevice, uploaded: NanoVDBWebGPUGrid, byteOffsets: readonly number[]): 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 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>;
${traversalWGSL}
@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 } },
] });
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();
return result;
}
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);
@@ -486,6 +590,7 @@ export async function renderNanoVDBFloat32WebGPU(
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;
@@ -501,11 +606,13 @@ export async function renderNanoVDBFloat32WebGPU(
? [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);
f32.set([material.densityScale, material.emissionScale, material.anisotropy, Math.max(0.01, gridDefinition.voxelSize[2])], 40);
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);
@@ -528,7 +635,7 @@ 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,
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,
@@ -539,9 +646,12 @@ struct Params {
@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(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}
@@ -553,39 +663,59 @@ fn blackbody_color(kelvin: f32) -> vec3<f32> {
@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));
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 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 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 z = params.index_min.z; z <= params.index_max.z; z += stride) {
var sample = sample_density(vec3<i32>(xy, z));
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(vec3<f32>(xy_position, f32(z) + 0.5));
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(vec3<i32>(xy, z));
if (params.interpolation == 1u) { color_sample = sample_color_linear(vec3<f32>(xy_position, f32(z) + 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(vec3<i32>(xy, z));
if (params.interpolation == 1u) { temperature_sample = temperature_sample_density_linear(vec3<f32>(xy_position, f32(z) + 0.5)); }
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(vec3<i32>(xy, z));
if (params.interpolation == 1u) { emission_sample = emission_sample_density_linear(vec3<f32>(xy_position, f32(z) + 0.5)); }
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;
@@ -600,9 +730,18 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>) {
{ 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 } });
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();