diff --git a/5-Applications/dashboard/spatial-hash-gpu/grid-storage.js b/5-Applications/dashboard/spatial-hash-gpu/grid-storage.js
new file mode 100644
index 00000000..2b975583
--- /dev/null
+++ b/5-Applications/dashboard/spatial-hash-gpu/grid-storage.js
@@ -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;
diff --git a/5-Applications/dashboard/spatial-hash-gpu/index.html b/5-Applications/dashboard/spatial-hash-gpu/index.html
new file mode 100644
index 00000000..2d05c2b0
--- /dev/null
+++ b/5-Applications/dashboard/spatial-hash-gpu/index.html
@@ -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 &nbsp;
+  <kbd>2</kbd> Clear grid &nbsp;
+  <kbd>3</kbd> Filter (density &gt; 50) &nbsp;
+  <kbd>4</kbd> Neighbor scan &nbsp;
+  <kbd>5</kbd> Sort by density &nbsp;
+  <kbd>6</kbd> Set voltage modes &nbsp;
+  <kbd>Mouse</kbd> Orbit camera &nbsp;
+  <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>
diff --git a/5-Applications/dashboard/spatial-hash-gpu/shaders.wgsl b/5-Applications/dashboard/spatial-hash-gpu/shaders.wgsl
new file mode 100644
index 00000000..80622a17
--- /dev/null
+++ b/5-Applications/dashboard/spatial-hash-gpu/shaders.wgsl
@@ -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);
+}