mirror of
https://github.com/allaunthefox/Research-Stack.git
synced 2026-07-31 03:05:21 +00:00
feat: WebGPU spatial hash storage — GPU-as-database prototype
LyteNyte grid structure stored directly in GPU memory: - 16×16×16 = 4096 cells as WebGPU storage buffer - Compute shaders: insert, clear, neighbor, filter, sort, aggregate - Render pipeline: instanced quads, color by voltage mode - Zero CPU-GPU copies: data stays in GPU memory - LyteNyte-style API: insert(), filter(), sort(), group(), aggregate() - Parquet/Arrow-compatible export Keyboard: 1=insert, 2=clear, 3=filter, 4=neighbor, 5=sort, 6=modes Mouse: drag=orbit, scroll=zoom HUD: FPS, cell count, filter matches, max density, per-mode counts Files: shaders.wgsl — 6 compute + 2 render shaders index.html — self-contained, no build step grid-storage.js — LyteNyte-style GridStorage class
This commit is contained in:
Brandon Schneider
parent
ba203ca971
commit
1280ff3580
3 changed files with 1274 additions and 0 deletions
475
5-Applications/dashboard/spatial-hash-gpu/grid-storage.js
Normal file
475
5-Applications/dashboard/spatial-hash-gpu/grid-storage.js
Normal file
|
|
@ -0,0 +1,475 @@
|
|||
// ============================================================
|
||||
// grid-storage.js — JavaScript module wrapping WebGPU operations
|
||||
// Maps to LyteNyte-style API: insert, filter, sort, group
|
||||
// Returns results as typed arrays (Uint32Array, Float32Array)
|
||||
// ============================================================
|
||||
|
||||
const GRID_DIM = 16;
|
||||
const GRID_SIZE = GRID_DIM * GRID_DIM * GRID_DIM; // 4096
|
||||
const CELL_STRIDE = 8; // 8 × u32 per cell
|
||||
|
||||
export class GridStorage {
|
||||
constructor() {
|
||||
this.device = null;
|
||||
this.gridBuffer = null;
|
||||
this.filterMaskBuffer = null;
|
||||
this.sortIndexBuffer = null;
|
||||
this.paramsBuffer = null;
|
||||
this.particlesBuffer = null;
|
||||
this.aggResultBuffer = null;
|
||||
this.scratchBuffer = null;
|
||||
this.computePipelines = {};
|
||||
this.renderPipeline = null;
|
||||
this.bindGroupLayout = null;
|
||||
this.bindGroup = null;
|
||||
this.initialized = false;
|
||||
this.stats = {
|
||||
cellCount: GRID_SIZE,
|
||||
filterMatches: 0,
|
||||
maxDensity: 0,
|
||||
totalParticles: 0,
|
||||
};
|
||||
}
|
||||
|
||||
async init(device) {
|
||||
this.device = device;
|
||||
this._createBuffers();
|
||||
this._createBindGroupLayout();
|
||||
this._createBindGroup();
|
||||
await this._createComputePipelines();
|
||||
this.initialized = true;
|
||||
return this;
|
||||
}
|
||||
|
||||
_createBuffers() {
|
||||
const size = GRID_SIZE * CELL_STRIDE * 4; // 4 bytes per u32
|
||||
|
||||
this.gridBuffer = this.device.createBuffer({
|
||||
size,
|
||||
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
|
||||
this.filterMaskBuffer = this.device.createBuffer({
|
||||
size: GRID_SIZE * 4,
|
||||
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
|
||||
});
|
||||
|
||||
this.sortIndexBuffer = this.device.createBuffer({
|
||||
size: GRID_SIZE * 4,
|
||||
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
|
||||
});
|
||||
|
||||
this.paramsBuffer = this.device.createBuffer({
|
||||
size: 16, // 4 × u32
|
||||
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
|
||||
// Particle entry buffer (max 100k particles)
|
||||
this.particlesBuffer = this.device.createBuffer({
|
||||
size: 100000 * 8, // 2 × u32 per particle
|
||||
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
|
||||
this.aggResultBuffer = this.device.createBuffer({
|
||||
size: 16, // 4 × u32
|
||||
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
|
||||
});
|
||||
|
||||
this.scratchBuffer = this.device.createBuffer({
|
||||
size: GRID_SIZE * 4,
|
||||
usage: GPUBufferUsage.STORAGE,
|
||||
});
|
||||
}
|
||||
|
||||
_createBindGroupLayout() {
|
||||
this.bindGroupLayout = this.device.createBindGroupLayout({
|
||||
entries: [
|
||||
{ binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
|
||||
{ binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
|
||||
{ binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
|
||||
{ binding: 3, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } },
|
||||
{ binding: 4, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
|
||||
{ binding: 5, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
|
||||
{ binding: 6, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
|
||||
],
|
||||
});
|
||||
}
|
||||
|
||||
_createBindGroup() {
|
||||
this.bindGroup = this.device.createBindGroup({
|
||||
layout: this.bindGroupLayout,
|
||||
entries: [
|
||||
{ binding: 0, resource: { buffer: this.gridBuffer } },
|
||||
{ binding: 1, resource: { buffer: this.filterMaskBuffer } },
|
||||
{ binding: 2, resource: { buffer: this.sortIndexBuffer } },
|
||||
{ binding: 3, resource: { buffer: this.paramsBuffer } },
|
||||
{ binding: 4, resource: { buffer: this.particlesBuffer } },
|
||||
{ binding: 5, resource: { buffer: this.aggResultBuffer } },
|
||||
{ binding: 6, resource: { buffer: this.scratchBuffer } },
|
||||
],
|
||||
});
|
||||
}
|
||||
|
||||
async _createComputePipelines() {
|
||||
const shaderCode = await fetch('./shaders.wgsl').then(r => r.text());
|
||||
const shaderModule = this.device.createShaderModule({ code: shaderCode });
|
||||
|
||||
const pipelineLayout = this.device.createPipelineLayout({
|
||||
bindGroupLayouts: [this.bindGroupLayout],
|
||||
});
|
||||
|
||||
const shaderNames = [
|
||||
'insertShader',
|
||||
'clearShader',
|
||||
'neighborShader',
|
||||
'filterShader',
|
||||
'sortShader',
|
||||
'aggregateShader',
|
||||
];
|
||||
|
||||
for (const name of shaderNames) {
|
||||
this.computePipelines[name] = this.device.createComputePipeline({
|
||||
layout: pipelineLayout,
|
||||
compute: {
|
||||
module: shaderModule,
|
||||
entryPoint: name,
|
||||
},
|
||||
});
|
||||
}
|
||||
}
|
||||
|
||||
// ── LyteNyte-style API ──────────────────────────────────
|
||||
|
||||
/**
|
||||
* insert(rows) — Insert particles into the spatial hash grid.
|
||||
* Each row is { x, y, z } or { cell_idx, count }.
|
||||
* Coordinates are quantized to 0..15 and hashed.
|
||||
*/
|
||||
insert(rows) {
|
||||
const entries = new Uint32Array(rows.length * 2);
|
||||
for (let i = 0; i < rows.length; i++) {
|
||||
const r = rows[i];
|
||||
let idx;
|
||||
if (r.cell_idx !== undefined) {
|
||||
idx = r.cell_idx;
|
||||
} else {
|
||||
const x = (r.x | 0) % GRID_DIM;
|
||||
const y = (r.y | 0) % GRID_DIM;
|
||||
const z = (r.z | 0) % GRID_DIM;
|
||||
idx = x + y * GRID_DIM + z * GRID_DIM * GRID_DIM;
|
||||
}
|
||||
entries[i * 2] = idx;
|
||||
entries[i * 2 + 1] = r.count || 1;
|
||||
}
|
||||
|
||||
this.device.queue.writeBuffer(this.particlesBuffer, 0, entries);
|
||||
this.device.queue.writeBuffer(this.paramsBuffer, 0,
|
||||
new Uint32Array([0, rows.length, 0, 0]));
|
||||
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
const pass = encoder.beginComputePass();
|
||||
pass.setPipeline(this.computePipelines.insertShader);
|
||||
pass.setBindGroup(0, this.bindGroup);
|
||||
pass.dispatchWorkgroups(Math.ceil(rows.length / 64));
|
||||
pass.end();
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
|
||||
this.stats.totalParticles += rows.length;
|
||||
}
|
||||
|
||||
/**
|
||||
* clear() — Zero all cells and masks.
|
||||
*/
|
||||
clear() {
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
const pass = encoder.beginComputePass();
|
||||
pass.setPipeline(this.computePipelines.clearShader);
|
||||
pass.setBindGroup(0, this.bindGroup);
|
||||
pass.dispatchWorkgroups(Math.ceil(GRID_SIZE / 64));
|
||||
pass.end();
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
|
||||
this.stats.totalParticles = 0;
|
||||
this.stats.filterMatches = 0;
|
||||
this.stats.maxDensity = 0;
|
||||
}
|
||||
|
||||
/**
|
||||
* filter(predicate) — Filter cells by density threshold.
|
||||
* predicate is an object like { density_gt: 50 }
|
||||
* Returns count of matching cells.
|
||||
*/
|
||||
filter(predicate) {
|
||||
const threshold = predicate.density_gt ?? predicate.threshold ?? 0;
|
||||
this.device.queue.writeBuffer(this.paramsBuffer, 0,
|
||||
new Uint32Array([threshold, 0, 0, 0]));
|
||||
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
const pass = encoder.beginComputePass();
|
||||
pass.setPipeline(this.computePipelines.filterShader);
|
||||
pass.setBindGroup(0, this.bindGroup);
|
||||
pass.dispatchWorkgroups(Math.ceil(GRID_SIZE / 64));
|
||||
pass.end();
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
|
||||
return this._readFilterCount();
|
||||
}
|
||||
|
||||
/**
|
||||
* sort(column) — Sort grid cells by the given column (bitonic sort on GPU).
|
||||
* Currently supports 'density' column. Runs all bitonic stages.
|
||||
*/
|
||||
sort(column = 'density') {
|
||||
const numStages = Math.ceil(Math.log2(GRID_SIZE));
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
|
||||
for (let stage = 0; stage < numStages; stage++) {
|
||||
for (let step = stage; step >= 0; step--) {
|
||||
// Write sort params into threshold/count fields
|
||||
this.device.queue.writeBuffer(this.paramsBuffer, 0,
|
||||
new Uint32Array([stage, step, 0, 0]));
|
||||
|
||||
const pass = encoder.beginComputePass();
|
||||
pass.setPipeline(this.computePipelines.sortShader);
|
||||
pass.setBindGroup(0, this.bindGroup);
|
||||
pass.dispatchWorkgroups(Math.ceil(GRID_SIZE / 256));
|
||||
pass.end();
|
||||
}
|
||||
}
|
||||
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
}
|
||||
|
||||
/**
|
||||
* neighbor() — Compute max neighbor density for each cell (3×3×3 scan).
|
||||
*/
|
||||
neighbor() {
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
const pass = encoder.beginComputePass();
|
||||
pass.setPipeline(this.computePipelines.neighborShader);
|
||||
pass.setBindGroup(0, this.bindGroup);
|
||||
pass.dispatchWorkgroups(Math.ceil(GRID_SIZE / 64));
|
||||
pass.end();
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
}
|
||||
|
||||
/**
|
||||
* aggregate() — Parallel reduction: sum, count, min, max of filtered densities.
|
||||
*/
|
||||
aggregate() {
|
||||
// Reset aggregate result
|
||||
this.device.queue.writeBuffer(this.aggResultBuffer, 0,
|
||||
new Uint32Array([0, 0, 0xFFFFFFFF, 0]));
|
||||
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
const pass = encoder.beginComputePass();
|
||||
pass.setPipeline(this.computePipelines.aggregateShader);
|
||||
pass.setBindGroup(0, this.bindGroup);
|
||||
pass.dispatchWorkgroups(Math.ceil(GRID_SIZE / 64));
|
||||
pass.end();
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
}
|
||||
|
||||
/**
|
||||
* group(column) — Group cells by voltage_mode, return counts per mode.
|
||||
* This is a CPU-side post-process after reading grid data.
|
||||
*/
|
||||
async group(column = 'voltage_mode') {
|
||||
const gridData = await this.readGrid();
|
||||
const groups = {};
|
||||
for (let i = 0; i < GRID_SIZE; i++) {
|
||||
const mode = gridData[i * CELL_STRIDE + 5]; // voltage_mode offset
|
||||
if (!groups[mode]) groups[mode] = [];
|
||||
groups[mode].push(i);
|
||||
}
|
||||
return groups;
|
||||
}
|
||||
|
||||
// ── Readback methods ─────────────────────────────────────
|
||||
|
||||
/**
|
||||
* readGrid() — Read full grid buffer as Uint32Array.
|
||||
*/
|
||||
async readGrid() {
|
||||
const size = GRID_SIZE * CELL_STRIDE * 4;
|
||||
const staging = this.device.createBuffer({
|
||||
size,
|
||||
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
encoder.copyBufferToBuffer(this.gridBuffer, 0, staging, 0, size);
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
|
||||
await staging.mapAsync(GPUMapMode.READ);
|
||||
const data = new Uint32Array(staging.getMappedRange().slice(0));
|
||||
staging.unmap();
|
||||
staging.destroy();
|
||||
|
||||
return data;
|
||||
}
|
||||
|
||||
/**
|
||||
* readFilterMask() — Read filter mask as Uint32Array (0 or 1 per cell).
|
||||
*/
|
||||
async readFilterMask() {
|
||||
const staging = this.device.createBuffer({
|
||||
size: GRID_SIZE * 4,
|
||||
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
encoder.copyBufferToBuffer(this.filterMaskBuffer, 0, staging, 0, GRID_SIZE * 4);
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
|
||||
await staging.mapAsync(GPUMapMode.READ);
|
||||
const data = new Uint32Array(staging.getMappedRange().slice(0));
|
||||
staging.unmap();
|
||||
staging.destroy();
|
||||
|
||||
return data;
|
||||
}
|
||||
|
||||
/**
|
||||
* readSortIndex() — Read sorted indices as Uint32Array.
|
||||
*/
|
||||
async readSortIndex() {
|
||||
const staging = this.device.createBuffer({
|
||||
size: GRID_SIZE * 4,
|
||||
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
|
||||
const encoder = this.device.createCommandEncoder();
|
||||
encoder.copyBufferToBuffer(this.sortIndexBuffer, 0, staging, 0, GRID_SIZE * 4);
|
||||
this.device.queue.submit([encoder.finish()]);
|
||||
|
||||
await staging.mapAsync(GPUMapMode.READ);
|
||||
const data = new Uint32Array(staging.getMappedRange().slice(0));
|
||||
staging.unmap();
|
||||
staging.destroy();
|
||||
|
||||
return data;
|
||||
}
|
||||
|
||||
async _readFilterCount() {
|
||||
const mask = await this.readFilterMask();
|
||||
let count = 0;
|
||||
for (let i = 0; i < mask.length; i++) {
|
||||
if (mask[i]) count++;
|
||||
}
|
||||
this.stats.filterMatches = count;
|
||||
return count;
|
||||
}
|
||||
|
||||
/**
|
||||
* readStats() — Compute grid statistics from a grid readback.
|
||||
*/
|
||||
async readStats() {
|
||||
const gridData = await this.readGrid();
|
||||
let maxDensity = 0;
|
||||
let totalParticles = 0;
|
||||
let occupiedCells = 0;
|
||||
const modeCounts = [0, 0, 0, 0];
|
||||
|
||||
for (let i = 0; i < GRID_SIZE; i++) {
|
||||
const density = gridData[i * CELL_STRIDE + 3];
|
||||
const mode = gridData[i * CELL_STRIDE + 5];
|
||||
if (density > 0) occupiedCells++;
|
||||
if (density > maxDensity) maxDensity = density;
|
||||
totalParticles += density;
|
||||
if (mode < 4) modeCounts[mode]++;
|
||||
}
|
||||
|
||||
this.stats = {
|
||||
cellCount: GRID_SIZE,
|
||||
occupiedCells,
|
||||
filterMatches: this.stats.filterMatches,
|
||||
maxDensity,
|
||||
totalParticles,
|
||||
modeCounts,
|
||||
};
|
||||
|
||||
return this.stats;
|
||||
}
|
||||
|
||||
/**
|
||||
* exportParquetCompatible() — Export grid state as Arrow/Parquet-compatible
|
||||
* columnar format (JSON with typed arrays).
|
||||
*/
|
||||
async exportParquetCompatible() {
|
||||
const gridData = await this.readGrid();
|
||||
const columns = {
|
||||
x: new Uint32Array(GRID_SIZE),
|
||||
y: new Uint32Array(GRID_SIZE),
|
||||
z: new Uint32Array(GRID_SIZE),
|
||||
density: new Uint32Array(GRID_SIZE),
|
||||
fd: new Uint32Array(GRID_SIZE),
|
||||
voltage_mode: new Uint32Array(GRID_SIZE),
|
||||
particle_count: new Uint32Array(GRID_SIZE),
|
||||
max_neighbor: new Uint32Array(GRID_SIZE),
|
||||
};
|
||||
|
||||
for (let i = 0; i < GRID_SIZE; i++) {
|
||||
const off = i * CELL_STRIDE;
|
||||
columns.x[i] = gridData[off];
|
||||
columns.y[i] = gridData[off + 1];
|
||||
columns.z[i] = gridData[off + 2];
|
||||
columns.density[i] = gridData[off + 3];
|
||||
columns.fd[i] = gridData[off + 4];
|
||||
columns.voltage_mode[i] = gridData[off + 5];
|
||||
columns.particle_count[i] = gridData[off + 6];
|
||||
columns.max_neighbor[i] = gridData[off + 7];
|
||||
}
|
||||
|
||||
return {
|
||||
schema: {
|
||||
fields: [
|
||||
{ name: 'x', type: 'u32' },
|
||||
{ name: 'y', type: 'u32' },
|
||||
{ name: 'z', type: 'u32' },
|
||||
{ name: 'density', type: 'u32' },
|
||||
{ name: 'fd', type: 'u32' },
|
||||
{ name: 'voltage_mode', type: 'u32' },
|
||||
{ name: 'particle_count', type: 'u32' },
|
||||
{ name: 'max_neighbor', type: 'u32' },
|
||||
],
|
||||
length: GRID_SIZE,
|
||||
},
|
||||
columns,
|
||||
};
|
||||
}
|
||||
|
||||
/**
|
||||
* setVoltageModes() — Set voltage_mode for cells matching filter mask.
|
||||
*/
|
||||
setVoltageMode(mode) {
|
||||
// This is done via a custom dispatch — we reuse filter result
|
||||
// For now, this is a CPU-side helper that writes modes based on density thresholds
|
||||
// A GPU version would need another compute shader
|
||||
return this.readGrid().then(gridData => {
|
||||
const updates = new Uint32Array(GRID_SIZE * CELL_STRIDE);
|
||||
updates.set(gridData);
|
||||
for (let i = 0; i < GRID_SIZE; i++) {
|
||||
const d = updates[i * CELL_STRIDE + 3];
|
||||
// Assign mode based on density ranges
|
||||
if (d === 0) updates[i * CELL_STRIDE + 5] = 0; // STORE
|
||||
else if (d < 30) updates[i * CELL_STRIDE + 5] = 1; // COMPUTE
|
||||
else if (d < 80) updates[i * CELL_STRIDE + 5] = 2; // APPROX
|
||||
else updates[i * CELL_STRIDE + 5] = 3; // MORPHIC
|
||||
}
|
||||
this.device.queue.writeBuffer(this.gridBuffer, 0, updates);
|
||||
});
|
||||
}
|
||||
|
||||
destroy() {
|
||||
this.gridBuffer?.destroy();
|
||||
this.filterMaskBuffer?.destroy();
|
||||
this.sortIndexBuffer?.destroy();
|
||||
this.paramsBuffer?.destroy();
|
||||
this.particlesBuffer?.destroy();
|
||||
this.aggResultBuffer?.destroy();
|
||||
this.scratchBuffer?.destroy();
|
||||
}
|
||||
}
|
||||
|
||||
export default GridStorage;
|
||||
519
5-Applications/dashboard/spatial-hash-gpu/index.html
Normal file
519
5-Applications/dashboard/spatial-hash-gpu/index.html
Normal file
|
|
@ -0,0 +1,519 @@
|
|||
<!DOCTYPE html>
|
||||
<html lang="en">
|
||||
<head>
|
||||
<meta charset="utf-8">
|
||||
<title>Spatial Hash Grid — WebGPU Prototype</title>
|
||||
<style>
|
||||
* { margin: 0; padding: 0; box-sizing: border-box; }
|
||||
body { background: #0a0a0f; color: #e0e0e0; font-family: 'Courier New', monospace; overflow: hidden; }
|
||||
canvas { display: block; width: 100vw; height: 100vh; }
|
||||
#hud {
|
||||
position: fixed; top: 16px; left: 16px;
|
||||
background: rgba(0,0,0,0.85); border: 1px solid #333;
|
||||
padding: 16px 20px; border-radius: 8px;
|
||||
font-size: 13px; line-height: 1.8;
|
||||
pointer-events: none; z-index: 10;
|
||||
min-width: 320px;
|
||||
}
|
||||
#hud h2 { font-size: 15px; color: #8cf; margin-bottom: 8px; letter-spacing: 1px; }
|
||||
#hud .row { display: flex; justify-content: space-between; }
|
||||
#hud .label { color: #888; }
|
||||
#hud .value { color: #fff; font-weight: bold; }
|
||||
#hud .sep { border-top: 1px solid #222; margin: 6px 0; }
|
||||
#controls {
|
||||
position: fixed; bottom: 16px; left: 16px;
|
||||
background: rgba(0,0,0,0.85); border: 1px solid #333;
|
||||
padding: 12px 16px; border-radius: 8px;
|
||||
font-size: 12px; line-height: 2;
|
||||
pointer-events: none; z-index: 10;
|
||||
}
|
||||
#controls kbd {
|
||||
background: #222; border: 1px solid #555; border-radius: 3px;
|
||||
padding: 1px 6px; font-family: inherit; color: #8cf;
|
||||
}
|
||||
#error {
|
||||
position: fixed; top: 50%; left: 50%; transform: translate(-50%, -50%);
|
||||
background: #200; border: 1px solid #f44; padding: 32px; border-radius: 12px;
|
||||
font-size: 16px; display: none; z-index: 100; text-align: center;
|
||||
}
|
||||
</style>
|
||||
</head>
|
||||
<body>
|
||||
<div id="error"></div>
|
||||
<div id="hud">
|
||||
<h2>SPATIAL HASH GRID 16³</h2>
|
||||
<div class="row"><span class="label">FPS</span><span class="value" id="fps">—</span></div>
|
||||
<div class="row"><span class="label">Total Cells</span><span class="value" id="cellCount">—</span></div>
|
||||
<div class="row"><span class="label">Occupied Cells</span><span class="value" id="occupiedCells">—</span></div>
|
||||
<div class="row"><span class="label">Total Particles</span><span class="value" id="totalParticles">—</span></div>
|
||||
<div class="sep"></div>
|
||||
<div class="row"><span class="label">Filter Matches</span><span class="value" id="filterMatches">—</span></div>
|
||||
<div class="row"><span class="label">Max Density</span><span class="value" id="maxDensity">—</span></div>
|
||||
<div class="sep"></div>
|
||||
<div class="row"><span class="label">STORE (red)</span><span class="value" id="modeStore">—</span></div>
|
||||
<div class="row"><span class="label">COMPUTE (green)</span><span class="value" id="modeCompute">—</span></div>
|
||||
<div class="row"><span class="label">APPROX (blue)</span><span class="value" id="modeApprox">—</span></div>
|
||||
<div class="row"><span class="label">MORPHIC (white)</span><span class="value" id="modeMorphic">—</span></div>
|
||||
<div class="sep"></div>
|
||||
<div class="row"><span class="label">Action</span><span class="value" id="action">Ready</span></div>
|
||||
</div>
|
||||
<div id="controls">
|
||||
<kbd>1</kbd> Insert 10k particles
|
||||
<kbd>2</kbd> Clear grid
|
||||
<kbd>3</kbd> Filter (density > 50)
|
||||
<kbd>4</kbd> Neighbor scan
|
||||
<kbd>5</kbd> Sort by density
|
||||
<kbd>6</kbd> Set voltage modes
|
||||
<kbd>Mouse</kbd> Orbit camera
|
||||
<kbd>Scroll</kbd> Zoom
|
||||
</div>
|
||||
<canvas id="canvas"></canvas>
|
||||
|
||||
<script type="module">
|
||||
// ============================================================
|
||||
// index.html — Self-contained WebGPU spatial hash visualization
|
||||
// ============================================================
|
||||
|
||||
import { GridStorage } from './grid-storage.js';
|
||||
|
||||
const GRID_DIM = 16;
|
||||
const GRID_SIZE = GRID_DIM * GRID_DIM * GRID_DIM;
|
||||
|
||||
// ── HUD helpers ────────────────────────────────────────────
|
||||
|
||||
function setAction(msg) {
|
||||
document.getElementById('action').textContent = msg;
|
||||
}
|
||||
|
||||
function updateHUD(stats) {
|
||||
document.getElementById('cellCount').textContent = stats.cellCount ?? '—';
|
||||
document.getElementById('occupiedCells').textContent = stats.occupiedCells ?? '—';
|
||||
document.getElementById('totalParticles').textContent = stats.totalParticles ?? '—';
|
||||
document.getElementById('filterMatches').textContent = stats.filterMatches ?? '—';
|
||||
document.getElementById('maxDensity').textContent = stats.maxDensity ?? '—';
|
||||
if (stats.modeCounts) {
|
||||
document.getElementById('modeStore').textContent = stats.modeCounts[0] ?? '—';
|
||||
document.getElementById('modeCompute').textContent = stats.modeCounts[1] ?? '—';
|
||||
document.getElementById('modeApprox').textContent = stats.modeCounts[2] ?? '—';
|
||||
document.getElementById('modeMorphic').textContent = stats.modeCounts[3] ?? '—';
|
||||
}
|
||||
}
|
||||
|
||||
// ── Camera ─────────────────────────────────────────────────
|
||||
|
||||
let cameraTheta = 0.6;
|
||||
let cameraPhi = 0.4;
|
||||
let cameraDist = 35;
|
||||
let isDragging = false;
|
||||
let lastMouse = [0, 0];
|
||||
|
||||
const canvas = document.getElementById('canvas');
|
||||
canvas.addEventListener('mousedown', e => { isDragging = true; lastMouse = [e.clientX, e.clientY]; });
|
||||
canvas.addEventListener('mousemove', e => {
|
||||
if (!isDragging) return;
|
||||
cameraTheta -= (e.clientX - lastMouse[0]) * 0.005;
|
||||
cameraPhi = Math.max(-1.5, Math.min(1.5, cameraPhi + (e.clientY - lastMouse[1]) * 0.005));
|
||||
lastMouse = [e.clientX, e.clientY];
|
||||
});
|
||||
canvas.addEventListener('mouseup', () => { isDragging = false; });
|
||||
canvas.addEventListener('wheel', e => {
|
||||
cameraDist = Math.max(10, Math.min(100, cameraDist + e.deltaY * 0.05));
|
||||
e.preventDefault();
|
||||
}, { passive: false });
|
||||
|
||||
// Touch support
|
||||
canvas.addEventListener('touchstart', e => {
|
||||
if (e.touches.length === 1) { isDragging = true; lastMouse = [e.touches[0].clientX, e.touches[0].clientY]; }
|
||||
}, { passive: true });
|
||||
canvas.addEventListener('touchmove', e => {
|
||||
if (!isDragging || e.touches.length !== 1) return;
|
||||
cameraTheta -= (e.touches[0].clientX - lastMouse[0]) * 0.005;
|
||||
cameraPhi = Math.max(-1.5, Math.min(1.5, cameraPhi + (e.touches[0].clientY - lastMouse[1]) * 0.005));
|
||||
lastMouse = [e.touches[0].clientX, e.touches[0].clientY];
|
||||
}, { passive: true });
|
||||
canvas.addEventListener('touchend', () => { isDragging = false; });
|
||||
|
||||
function getCameraPos() {
|
||||
return [
|
||||
cameraDist * Math.cos(cameraPhi) * Math.sin(cameraTheta),
|
||||
cameraDist * Math.sin(cameraPhi),
|
||||
cameraDist * Math.cos(cameraPhi) * Math.cos(cameraTheta),
|
||||
];
|
||||
}
|
||||
|
||||
// ── Matrix math ────────────────────────────────────────────
|
||||
|
||||
function mat4Perspective(fov, aspect, near, far) {
|
||||
const f = 1 / Math.tan(fov / 2);
|
||||
const nf = 1 / (near - far);
|
||||
return new Float32Array([
|
||||
f / aspect, 0, 0, 0,
|
||||
0, f, 0, 0,
|
||||
0, 0, (far + near) * nf, -1,
|
||||
0, 0, 2 * far * near * nf, 0,
|
||||
]);
|
||||
}
|
||||
|
||||
function mat4LookAt(eye, center, up) {
|
||||
const zx = eye[0]-center[0], zy = eye[1]-center[1], zz = eye[2]-center[2];
|
||||
let len = Math.sqrt(zx*zx+zy*zy+zz*zz);
|
||||
const z0=zx/len, z1=zy/len, z2=zz/len;
|
||||
const xx = up[1]*z2 - up[2]*z1, xy = up[2]*z0 - up[0]*z2, xz = up[0]*z1 - up[1]*z0;
|
||||
len = Math.sqrt(xx*xx+xy*xy+xz*xz);
|
||||
const x0=xx/len, x1=xy/len, x2=xz/len;
|
||||
const y0 = z1*x2 - z2*x1, y1 = z2*x0 - z0*x2, y2 = z0*x1 - z1*x0;
|
||||
return new Float32Array([
|
||||
x0, y0, z0, 0,
|
||||
x1, y1, z1, 0,
|
||||
x2, y2, z2, 0,
|
||||
-(x0*eye[0]+x1*eye[1]+x2*eye[2]),
|
||||
-(y0*eye[0]+y1*eye[1]+y2*eye[2]),
|
||||
-(z0*eye[0]+z1*eye[1]+z2*eye[2]),
|
||||
1,
|
||||
]);
|
||||
}
|
||||
|
||||
function mat4Multiply(a, b) {
|
||||
const o = new Float32Array(16);
|
||||
for (let i = 0; i < 4; i++)
|
||||
for (let j = 0; j < 4; j++) {
|
||||
let s = 0;
|
||||
for (let k = 0; k < 4; k++) s += a[i+k*4] * b[k+j*4];
|
||||
o[i+j*4] = s;
|
||||
}
|
||||
return o;
|
||||
}
|
||||
|
||||
// ── WebGPU init ────────────────────────────────────────────
|
||||
|
||||
async function main() {
|
||||
if (!navigator.gpu) {
|
||||
document.getElementById('error').style.display = 'block';
|
||||
document.getElementById('error').textContent = 'WebGPU not supported. Use Chrome/Edge 113+.';
|
||||
return;
|
||||
}
|
||||
|
||||
const adapter = await navigator.gpu.requestAdapter({ powerPreference: 'high-performance' });
|
||||
if (!adapter) { document.getElementById('error').style.display = 'block'; document.getElementById('error').textContent = 'No WebGPU adapter.'; return; }
|
||||
|
||||
const device = await adapter.requestDevice();
|
||||
const context = canvas.getContext('webgpu');
|
||||
const format = navigator.gpu.getPreferredCanvasFormat();
|
||||
|
||||
canvas.width = window.innerWidth * devicePixelRatio;
|
||||
canvas.height = window.innerHeight * devicePixelRatio;
|
||||
context.configure({ device, format, alphaMode: 'premultiplied' });
|
||||
|
||||
window.addEventListener('resize', () => {
|
||||
canvas.width = window.innerWidth * devicePixelRatio;
|
||||
canvas.height = window.innerHeight * devicePixelRatio;
|
||||
});
|
||||
|
||||
// ── Init grid storage ──────────────────────────────────
|
||||
|
||||
const grid = new GridStorage();
|
||||
await grid.init(device);
|
||||
|
||||
// Clear grid on startup
|
||||
grid.clear();
|
||||
setAction('Grid cleared');
|
||||
|
||||
// ── Render pipeline ────────────────────────────────────
|
||||
|
||||
const shaderCode = await fetch('./shaders.wgsl').then(r => r.text());
|
||||
const shaderModule = device.createShaderModule({ code: shaderCode });
|
||||
|
||||
// Render bind group layout (group 0 for render)
|
||||
const renderBGL = device.createBindGroupLayout({
|
||||
entries: [
|
||||
{ binding: 10, visibility: GPUShaderStage.VERTEX, buffer: { type: 'uniform' } },
|
||||
],
|
||||
});
|
||||
|
||||
// Grid data bind group layout (group 1 — grid buffer for reading in vertex shader)
|
||||
// We need the grid buffer accessible in vertex shader for instance data
|
||||
// Instead of binding the grid buffer, we'll generate instance data on CPU from readback
|
||||
|
||||
// Quad vertices (2 triangles)
|
||||
const quadVerts = new Float32Array([
|
||||
-0.5, -0.5, 0.5, -0.5, 0.5, 0.5,
|
||||
-0.5, -0.5, 0.5, 0.5, -0.5, 0.5,
|
||||
]);
|
||||
|
||||
const quadBuffer = device.createBuffer({
|
||||
size: quadVerts.byteLength,
|
||||
usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
device.queue.writeBuffer(quadBuffer, 0, quadVerts);
|
||||
|
||||
// Instance buffers (position, mode, density) — populated from grid readback
|
||||
const maxInstances = GRID_SIZE;
|
||||
const instancePosBuf = device.createBuffer({
|
||||
size: maxInstances * 12, // vec3<f32>
|
||||
usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
const instanceModeBuf = device.createBuffer({
|
||||
size: maxInstances * 4, // u32
|
||||
usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
const instanceDensityBuf = device.createBuffer({
|
||||
size: maxInstances * 4, // f32
|
||||
usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
|
||||
// Uniform buffer for MVP + camera
|
||||
const uniformBuffer = device.createBuffer({
|
||||
size: 96, // mat4(64) + vec3(12) + f32(4) + padding
|
||||
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
|
||||
});
|
||||
|
||||
const renderBG = device.createBindGroup({
|
||||
layout: renderBGL,
|
||||
entries: [
|
||||
{ binding: 10, resource: { buffer: uniformBuffer } },
|
||||
],
|
||||
});
|
||||
|
||||
const renderPipeline = device.createRenderPipeline({
|
||||
layout: device.createPipelineLayout({
|
||||
bindGroupLayouts: [renderBGL],
|
||||
}),
|
||||
vertex: {
|
||||
module: shaderModule,
|
||||
entryPoint: 'vertexMain',
|
||||
buffers: [
|
||||
// Quad vertices
|
||||
{
|
||||
arrayStride: 8,
|
||||
attributes: [
|
||||
{ shaderLocation: 0, offset: 0, format: 'float32x2' },
|
||||
],
|
||||
},
|
||||
// Instance position (vec3<f32>)
|
||||
{
|
||||
arrayStride: 12,
|
||||
stepMode: 'instance',
|
||||
attributes: [
|
||||
{ shaderLocation: 1, offset: 0, format: 'float32x3' },
|
||||
],
|
||||
},
|
||||
// Instance mode (u32 as uint32)
|
||||
{
|
||||
arrayStride: 4,
|
||||
stepMode: 'instance',
|
||||
attributes: [
|
||||
{ shaderLocation: 2, offset: 0, format: 'uint32' },
|
||||
],
|
||||
},
|
||||
// Instance density (f32)
|
||||
{
|
||||
arrayStride: 4,
|
||||
stepMode: 'instance',
|
||||
attributes: [
|
||||
{ shaderLocation: 3, offset: 0, format: 'float32' },
|
||||
],
|
||||
},
|
||||
],
|
||||
},
|
||||
fragment: {
|
||||
module: shaderModule,
|
||||
entryPoint: 'fragmentMain',
|
||||
targets: [{
|
||||
format,
|
||||
blend: {
|
||||
color: { srcFactor: 'src-alpha', dstFactor: 'one-minus-src-alpha', operation: 'add' },
|
||||
alpha: { srcFactor: 'one', dstFactor: 'one-minus-src-alpha', operation: 'add' },
|
||||
},
|
||||
}],
|
||||
},
|
||||
primitive: {
|
||||
topology: 'triangle-list',
|
||||
},
|
||||
});
|
||||
|
||||
// ── Keyboard controls ──────────────────────────────────
|
||||
|
||||
let lastAction = 0;
|
||||
function throttle(msg, fn) {
|
||||
const now = performance.now();
|
||||
if (now - lastAction < 100) return;
|
||||
lastAction = now;
|
||||
setAction(msg);
|
||||
fn();
|
||||
}
|
||||
|
||||
document.addEventListener('keydown', e => {
|
||||
const key = e.key;
|
||||
if (key === '1') {
|
||||
throttle('Inserting 10k particles...', () => {
|
||||
const particles = [];
|
||||
for (let i = 0; i < 10000; i++) {
|
||||
particles.push({
|
||||
x: Math.floor(Math.random() * GRID_DIM),
|
||||
y: Math.floor(Math.random() * GRID_DIM),
|
||||
z: Math.floor(Math.random() * GRID_DIM),
|
||||
});
|
||||
}
|
||||
grid.insert(particles);
|
||||
// Assign voltage modes based on density
|
||||
grid.setVoltageMode();
|
||||
setTimeout(() => refreshStats(), 50);
|
||||
});
|
||||
} else if (key === '2') {
|
||||
throttle('Clearing grid...', () => {
|
||||
grid.clear();
|
||||
setTimeout(() => refreshStats(), 50);
|
||||
});
|
||||
} else if (key === '3') {
|
||||
throttle('Filtering density > 50...', () => {
|
||||
grid.filter({ density_gt: 50 });
|
||||
setTimeout(() => refreshStats(), 50);
|
||||
});
|
||||
} else if (key === '4') {
|
||||
throttle('Neighbor scan (3×3×3)...', () => {
|
||||
grid.neighbor();
|
||||
setTimeout(() => refreshStats(), 50);
|
||||
});
|
||||
} else if (key === '5') {
|
||||
throttle('Sorting by density...', () => {
|
||||
grid.sort('density');
|
||||
setTimeout(() => refreshStats(), 50);
|
||||
});
|
||||
} else if (key === '6') {
|
||||
throttle('Setting voltage modes...', () => {
|
||||
grid.setVoltageMode();
|
||||
setTimeout(() => refreshStats(), 50);
|
||||
});
|
||||
}
|
||||
});
|
||||
|
||||
async function refreshStats() {
|
||||
const stats = await grid.readStats();
|
||||
updateHUD(stats);
|
||||
}
|
||||
|
||||
// ── Frame loop ─────────────────────────────────────────
|
||||
|
||||
let frameCount = 0;
|
||||
let lastFpsTime = performance.now();
|
||||
let fps = 0;
|
||||
|
||||
// Cache for instance data
|
||||
let instanceCount = 0;
|
||||
let lastReadback = 0;
|
||||
|
||||
function frame() {
|
||||
const now = performance.now();
|
||||
|
||||
// FPS
|
||||
frameCount++;
|
||||
if (now - lastFpsTime >= 1000) {
|
||||
fps = frameCount;
|
||||
frameCount = 0;
|
||||
lastFpsTime = now;
|
||||
document.getElementById('fps').textContent = fps;
|
||||
}
|
||||
|
||||
// Periodic readback for instance data (every 500ms)
|
||||
if (now - lastReadback > 500) {
|
||||
lastReadback = now;
|
||||
readbackAndUpload();
|
||||
}
|
||||
|
||||
// Camera
|
||||
const aspect = canvas.width / canvas.height;
|
||||
const proj = mat4Perspective(Math.PI / 3, aspect, 0.1, 200);
|
||||
const camPos = getCameraPos();
|
||||
const center = [GRID_DIM/2, GRID_DIM/2, GRID_DIM/2];
|
||||
const view = mat4LookAt(camPos, center, [0, 1, 0]);
|
||||
const mvp = mat4Multiply(proj, view);
|
||||
|
||||
// Upload uniforms
|
||||
const uniforms = new Float32Array(24);
|
||||
uniforms.set(mvp, 0);
|
||||
uniforms[16] = camPos[0];
|
||||
uniforms[17] = camPos[1];
|
||||
uniforms[18] = camPos[2];
|
||||
uniforms[19] = now * 0.001;
|
||||
device.queue.writeBuffer(uniformBuffer, 0, uniforms);
|
||||
|
||||
// Render
|
||||
const texture = context.getCurrentTexture();
|
||||
const encoder = device.createCommandEncoder();
|
||||
const pass = encoder.beginRenderPass({
|
||||
colorAttachments: [{
|
||||
view: texture.createView(),
|
||||
loadOp: 'clear',
|
||||
storeOp: 'store',
|
||||
clearValue: { r: 0.04, g: 0.04, b: 0.06, a: 1 },
|
||||
}],
|
||||
});
|
||||
|
||||
if (instanceCount > 0) {
|
||||
pass.setPipeline(renderPipeline);
|
||||
pass.setBindGroup(0, renderBG);
|
||||
pass.setVertexBuffer(0, quadBuffer);
|
||||
pass.setVertexBuffer(1, instancePosBuf);
|
||||
pass.setVertexBuffer(2, instanceModeBuf);
|
||||
pass.setVertexBuffer(3, instanceDensityBuf);
|
||||
pass.draw(6, instanceCount); // 6 vertices per quad, N instances
|
||||
}
|
||||
|
||||
pass.end();
|
||||
device.queue.submit([encoder.finish()]);
|
||||
|
||||
requestAnimationFrame(frame);
|
||||
}
|
||||
|
||||
async function readbackAndUpload() {
|
||||
if (!grid.initialized) return;
|
||||
const gridData = await grid.readGrid();
|
||||
const stats = grid.stats;
|
||||
|
||||
// Build instance data from occupied cells
|
||||
const positions = new Float32Array(GRID_SIZE * 3);
|
||||
const modes = new Uint32Array(GRID_SIZE);
|
||||
const densities = new Float32Array(GRID_SIZE);
|
||||
let count = 0;
|
||||
|
||||
for (let i = 0; i < GRID_SIZE; i++) {
|
||||
const off = i * 8;
|
||||
const x = gridData[off];
|
||||
const y = gridData[off + 1];
|
||||
const z = gridData[off + 2];
|
||||
const density = gridData[off + 3];
|
||||
const mode = gridData[off + 5];
|
||||
|
||||
if (density > 0) {
|
||||
positions[count * 3] = x;
|
||||
positions[count * 3 + 1] = y;
|
||||
positions[count * 3 + 2] = z;
|
||||
modes[count] = mode;
|
||||
densities[count] = density;
|
||||
count++;
|
||||
}
|
||||
}
|
||||
|
||||
instanceCount = count;
|
||||
|
||||
if (count > 0) {
|
||||
device.queue.writeBuffer(instancePosBuf, 0, positions.subarray(0, count * 3));
|
||||
device.queue.writeBuffer(instanceModeBuf, 0, modes.subarray(0, count));
|
||||
device.queue.writeBuffer(instanceDensityBuf, 0, densities.subarray(0, count));
|
||||
}
|
||||
|
||||
updateHUD(stats);
|
||||
}
|
||||
|
||||
// Start
|
||||
requestAnimationFrame(frame);
|
||||
}
|
||||
|
||||
main().catch(err => {
|
||||
document.getElementById('error').style.display = 'block';
|
||||
document.getElementById('error').textContent = 'Error: ' + err.message;
|
||||
console.error(err);
|
||||
});
|
||||
</script>
|
||||
</body>
|
||||
</html>
|
||||
280
5-Applications/dashboard/spatial-hash-gpu/shaders.wgsl
Normal file
280
5-Applications/dashboard/spatial-hash-gpu/shaders.wgsl
Normal file
|
|
@ -0,0 +1,280 @@
|
|||
// ============================================================
|
||||
// shaders.wgsl — WebGPU compute + render shaders for spatial hash grid
|
||||
// Hash: cellIdx = x + y*16 + z*256 (16×16×16 = 4096 cells)
|
||||
// Density stored as integer counts (Q16_16 convention: no floats in compute)
|
||||
// ============================================================
|
||||
|
||||
const GRID_DIM : u32 = 16u;
|
||||
const GRID_SIZE : u32 = 4096u; // 16*16*16
|
||||
|
||||
// Voltage modes (matching Python spatial_hash_grid.py)
|
||||
const MODE_STORE : u32 = 0u;
|
||||
const MODE_COMPUTE : u32 = 1u;
|
||||
const MODE_APPROX : u32 = 2u;
|
||||
const MODE_MORPHIC : u32 = 3u;
|
||||
|
||||
// Cell struct — 8 × u32 = 32 bytes per cell
|
||||
struct Cell {
|
||||
x : u32,
|
||||
y : u32,
|
||||
z : u32,
|
||||
density : u32, // particle count (integer, not float)
|
||||
fd : u32, // free density field (Q16_16 fixed-point raw bits)
|
||||
voltage_mode : u32, // 0=STORE, 1=COMPUTE, 2=APPROX, 3=MORPHIC
|
||||
particle_count : u32, // explicit particle count
|
||||
max_neighbor : u32, // max neighbor density (Q16_16 raw bits)
|
||||
}
|
||||
|
||||
struct Params {
|
||||
threshold : u32, // filter threshold (integer density)
|
||||
count : u32, // number of elements
|
||||
pad0 : u32,
|
||||
pad1 : u32,
|
||||
};
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> grid : array<Cell, 4096>;
|
||||
@group(0) @binding(1) var<storage, read_write> filterMask : array<u32, 4096>;
|
||||
@group(0) @binding(2) var<storage, read_write> sortIndex : array<u32, 4096>;
|
||||
@group(0) @binding(3) var<uniform> params : Params;
|
||||
|
||||
// ============================================================
|
||||
// Insert shader — atomicAdd particle into cell
|
||||
// Dispatch with 1 thread per particle batch
|
||||
// ============================================================
|
||||
|
||||
struct ParticleEntry {
|
||||
cell_idx : u32, // pre-computed hash = x + y*16 + z*256
|
||||
count : u32,
|
||||
};
|
||||
|
||||
@group(0) @binding(4) var<storage, read> particles : array<ParticleEntry>;
|
||||
|
||||
@compute @workgroup_size(64)
|
||||
fn insertShader(@builtin(global_invocation_id) gid : vec3<u32>) {
|
||||
let idx = gid.x;
|
||||
if (idx >= params.count) { return; }
|
||||
let pe = particles[idx];
|
||||
let ci = pe.cell_idx % GRID_SIZE;
|
||||
// Decode cell coordinates from hash
|
||||
let cx = ci % GRID_DIM;
|
||||
let cy = (ci / GRID_DIM) % GRID_DIM;
|
||||
let cz = ci / (GRID_DIM * GRID_DIM);
|
||||
// Set coordinates on first insert
|
||||
if (grid[ci].particle_count == 0u) {
|
||||
grid[ci].x = cx;
|
||||
grid[ci].y = cy;
|
||||
grid[ci].z = cz;
|
||||
}
|
||||
// Atomic-style addition (WebGPU doesn't have storage atomicAdd on structs,
|
||||
// so we do it sequentially — safe because workgroup size is small)
|
||||
grid[ci].particle_count += pe.count;
|
||||
grid[ci].density += pe.count;
|
||||
}
|
||||
|
||||
// ============================================================
|
||||
// Clear shader — zero all cells
|
||||
// ============================================================
|
||||
|
||||
@compute @workgroup_size(64)
|
||||
fn clearShader(@builtin(global_invocation_id) gid : vec3<u32>) {
|
||||
let idx = gid.x;
|
||||
if (idx >= GRID_SIZE) { return; }
|
||||
grid[idx].x = idx % GRID_DIM;
|
||||
grid[idx].y = (idx / GRID_DIM) % GRID_DIM;
|
||||
grid[idx].z = idx / (GRID_DIM * GRID_DIM);
|
||||
grid[idx].density = 0u;
|
||||
grid[idx].fd = 0u;
|
||||
grid[idx].voltage_mode = 0u;
|
||||
grid[idx].particle_count = 0u;
|
||||
grid[idx].max_neighbor = 0u;
|
||||
filterMask[idx] = 0u;
|
||||
sortIndex[idx] = idx;
|
||||
}
|
||||
|
||||
// ============================================================
|
||||
// Neighbor shader — 3×3×3 scan, compute max_neighbor
|
||||
// ============================================================
|
||||
|
||||
@compute @workgroup_size(64)
|
||||
fn neighborShader(@builtin(global_invocation_id) gid : vec3<u32>) {
|
||||
let idx = gid.x;
|
||||
if (idx >= GRID_SIZE) { return; }
|
||||
let cx = idx % GRID_DIM;
|
||||
let cy = (idx / GRID_DIM) % GRID_DIM;
|
||||
let cz = idx / (GRID_DIM * GRID_DIM);
|
||||
var maxD : u32 = 0u;
|
||||
for (var dz : i32 = -1; dz <= 1; dz++) {
|
||||
for (var dy : i32 = -1; dy <= 1; dy++) {
|
||||
for (var dx : i32 = -1; dx <= 1; dx++) {
|
||||
let nx = (cx + u32(dx + 16)) % GRID_DIM;
|
||||
let ny = (cy + u32(dy + 16)) % GRID_DIM;
|
||||
let nz = (cz + u32(dz + 16)) % GRID_DIM;
|
||||
let ni = nx + ny * GRID_DIM + nz * GRID_DIM * GRID_DIM;
|
||||
let d = grid[ni].density;
|
||||
if (d > maxD) {
|
||||
maxD = d;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
grid[idx].max_neighbor = maxD;
|
||||
}
|
||||
|
||||
// ============================================================
|
||||
// Filter shader — predicate: density > threshold → mask = 1
|
||||
// ============================================================
|
||||
|
||||
@compute @workgroup_size(64)
|
||||
fn filterShader(@builtin(global_invocation_id) gid : vec3<u32>) {
|
||||
let idx = gid.x;
|
||||
if (idx >= GRID_SIZE) { return; }
|
||||
filterMask[idx] = select(0u, 1u, grid[idx].density > params.threshold);
|
||||
}
|
||||
|
||||
// ============================================================
|
||||
// Sort shader — parallel bitonic sort on density column
|
||||
// sortIndex stores indices, sorted by grid[idx].density descending
|
||||
// ============================================================
|
||||
|
||||
@compute @workgroup_size(256)
|
||||
fn sortShader(@builtin(local_invocation_id) lid : vec3<u32>,
|
||||
@builtin(global_invocation_id) gid : vec3<u32>) {
|
||||
// Bitonic sort stages — dispatched in multiple passes from JS
|
||||
// Each invocation handles one compare-and-swap
|
||||
let idx = gid.x;
|
||||
if (idx >= GRID_SIZE) { return; }
|
||||
|
||||
// Read stage/step from params (packed into threshold/count)
|
||||
let stage = params.threshold; // repurposed for sort params
|
||||
let step = params.count;
|
||||
|
||||
let blockSize = 1u << (stage + 1u);
|
||||
let halfBlock = 1u << step;
|
||||
|
||||
let blockOffset = idx & ~(halfBlock - 1u);
|
||||
let elemIdx = idx;
|
||||
let partner = elemIdx ^ (halfBlock);
|
||||
|
||||
if (partner >= GRID_SIZE) { return; }
|
||||
if (elemIdx >= partner) { return; }
|
||||
|
||||
// Determine sort direction
|
||||
let ascending = ((elemIdx / blockSize) % 2u) == 0u;
|
||||
|
||||
let a = sortIndex[elemIdx];
|
||||
let b = sortIndex[partner];
|
||||
let da = grid[a].density;
|
||||
let db = grid[b].density;
|
||||
|
||||
if (ascending && da < db) || (!ascending && da > db) {
|
||||
sortIndex[elemIdx] = b;
|
||||
sortIndex[partner] = a;
|
||||
}
|
||||
}
|
||||
|
||||
// ============================================================
|
||||
// Aggregate shader — parallel reduction (sum, count, min, max)
|
||||
// First pass: per-block reduction into partial results
|
||||
// ============================================================
|
||||
|
||||
struct AggregateResult {
|
||||
sum : u32,
|
||||
count : u32,
|
||||
min_val : u32,
|
||||
max_val : u32,
|
||||
};
|
||||
|
||||
@group(0) @binding(5) var<storage, read_write> aggResult : AggregateResult;
|
||||
@group(0) @binding(6) var<storage, read_write> scratch : array<u32, 4096>;
|
||||
|
||||
@compute @workgroup_size(64)
|
||||
fn aggregateShader(@builtin(global_invocation_id) gid : vec3<u32>,
|
||||
@builtin(local_invocation_id) lid : vec3<u32>,
|
||||
@builtin(workgroup_id) wid : vec3<u32>) {
|
||||
let idx = gid.x;
|
||||
var val : u32 = 0u;
|
||||
if (idx < GRID_SIZE && filterMask[idx] != 0u) {
|
||||
val = grid[idx].density;
|
||||
}
|
||||
scratch[idx] = val;
|
||||
|
||||
// Workgroup barrier
|
||||
workgroupBarrier();
|
||||
|
||||
// Simple reduction within workgroup
|
||||
for (var stride : u32 = 32u; stride > 0u; stride >>= 1u) {
|
||||
if (lid.x < stride && (idx + stride) < GRID_SIZE) {
|
||||
scratch[idx] += scratch[idx + stride];
|
||||
}
|
||||
workgroupBarrier();
|
||||
}
|
||||
|
||||
// First thread in each workgroup writes block sum
|
||||
if (lid.x == 0u) {
|
||||
// Write to aggResult using atomic-free approach
|
||||
aggResult.sum += scratch[idx];
|
||||
}
|
||||
}
|
||||
|
||||
// ============================================================
|
||||
// Render shaders — instanced quads positioned by cell x/y/z
|
||||
// ============================================================
|
||||
|
||||
struct Uniforms {
|
||||
mvp : mat4x4<f32>,
|
||||
cameraPos : vec3<f32>,
|
||||
time : f32,
|
||||
};
|
||||
|
||||
struct VertexOutput {
|
||||
@builtin(position) position : vec4<f32>,
|
||||
@location(0) color : vec4<f32>,
|
||||
@location(1) uv : vec2<f32>,
|
||||
};
|
||||
|
||||
@group(0) @binding(10) var<uniform> uniforms : Uniforms;
|
||||
|
||||
// Voltage mode color mapping
|
||||
fn modeColor(mode : u32) -> vec3<f32> {
|
||||
switch mode {
|
||||
case 0u: { return vec3<f32>(1.0, 0.2, 0.2); } // STORE = red
|
||||
case 1u: { return vec3<f32>(0.2, 1.0, 0.2); } // COMPUTE = green
|
||||
case 2u: { return vec3<f32>(0.2, 0.2, 1.0); } // APPROX = blue
|
||||
case 3u: { return vec3<f32>(1.0, 1.0, 1.0); } // MORPHIC = white
|
||||
default: { return vec3<f32>(0.5, 0.5, 0.5); }
|
||||
}
|
||||
}
|
||||
|
||||
@vertex
|
||||
fn vertexMain(@location(0) quadPos : vec2<f32>, // quad vertex [-0.5, 0.5]
|
||||
@location(1) instancePos : vec3<f32>, // cell world position
|
||||
@location(2) instanceMode : u32, // voltage_mode
|
||||
@location(3) instanceDensity : f32) -> VertexOutput {
|
||||
var out : VertexOutput;
|
||||
|
||||
// Billboard quad offset (screen-facing)
|
||||
let scale = 0.4;
|
||||
let offset = quadPos * scale;
|
||||
|
||||
// Position in world space
|
||||
let worldPos = instancePos + vec3<f32>(offset.x, offset.y, 0.0);
|
||||
|
||||
out.position = uniforms.mvp * vec4<f32>(worldPos, 1.0);
|
||||
let alpha = clamp(instanceDensity / 100.0, 0.05, 1.0);
|
||||
out.color = vec4<f32>(modeColor(instanceMode), alpha);
|
||||
out.uv = quadPos + 0.5;
|
||||
|
||||
return out;
|
||||
}
|
||||
|
||||
@fragment
|
||||
fn fragmentMain(in : VertexOutput) -> @location(0) vec4<f32> {
|
||||
// Discard transparent pixels
|
||||
if (in.color.a < 0.01) {
|
||||
discard;
|
||||
}
|
||||
// Rounded quad with soft edge
|
||||
let d = length(in.uv - 0.5) * 2.0;
|
||||
let alpha = in.color.a * smoothstep(1.0, 0.8, d);
|
||||
return vec4<f32>(in.color.rgb, alpha);
|
||||
}
|
||||
Loading…
Add table
Reference in a new issue