Smooth look on WebGPU: WGSL port fed by GPU inputs and CPU overrides (t88 phase 2c)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

docs/PROJECT_MAP.md

@@ -149,25 +149,27 @@ The frame loop applies speed and queued steps while keeping the per-frame tick l
 
 ## Drawing
 
-- **The pixel look:**
+`SET.renderer` selects WebGPU or Legacy for either look. `src/render/cell.js` holds the shared
-  - `drawWorld` (`src/render/world.js`) fills one pixel per grain;
+CPU rules; `objects.js` adds bars, flyers and blueprint ghosts. The GPU canvas sits behind the
-  - `diagonal-screws.js` derives diagonal and straight end bends from `endJoin`, sharing casing geometry between pixels, smooth shading tags, ghosts and previews;
+transparent game canvas, which draws the existing overlays (machines, molds, bot, rain, arrows).
-  - `computeLight` (`light.js`) adds coloured light;
-  - overlays (machines, molds, the bot, rain, arrows) are drawn on the 2D canvas afterwards
-    (`machines.js`, `overlays.js`).
-- **The smooth look** (`src/render/smooth.js`) takes the buffers `drawWorld` fills (colour and
-  a class per cell, the colour behind, trees, fields, light, glow) and shades them with WebGL2.
-  - The cell classes are set in `world.js` and read in the shader: 4 is a crisp built piece,
-    6 a wall tile, 7/8 a belt, 9–12 a screw conveyor, 13/14 a filter.
-  - On software GL it falls back to pixels, unless `window.__allowSoftwareGL` is set (the test
-    tools set it).
 
-The pixel look can use WebGPU (`src/render/gpu/`): whole-world textures receive visible
+- **WebGPU** (`src/render/gpu/`): `textures.js` uploads visible world rectangles, `colour.js`
-strided uploads, `colour.js` colours plain grains, and `draw.js` selects rare CPU override
+  colours plain grains, and `draw.js` selects rare CPU override blocks and submits the passes.
-blocks and submits colour, bloom and drawing. `cell.js` shares the original CPU cell rules
+  `light.js` computes lighting; the CPU `computeLight` is never called by this path.
-between Legacy/Smooth and those overrides; `objects.js` adds bars, flyers and blueprint ghosts.
+  - Pixel uses nearest upscaling and separable GPU bloom (`shaders.js`).
-The GPU canvas sits behind the transparent game canvas, which draws all existing overlays.
+  - Smooth adds class, back colour/sky flag, tree/geometry tags and occupancy field outputs to
-CPU lighting stays in `light.js`. `SET.renderer` selects WebGPU or Legacy; Smooth uses WebGL2.
+    colour compute. Only this look uploads the CPU back/tree/field overrides and presence mask.
+  - `smooth-shader.js` and `smooth-slope.js` literally port the legacy GLSL shapes to WGSL.
+    `smooth.js` creates the smooth pipeline, field mip passes and bindings. Field AO samples
+    level 2.2; glow retains the legacy nine-tap filter. WebGPU smooth always draws full resolution.
+- **Legacy**: `drawWorld` fills every cell on the CPU, and `render/light.js` computes lighting.
+  Pixel uses Canvas2D; smooth (`render/smooth.js`, `smooth-shaders.js`, `slope-belt-shader.js`)
+  shades the four inputs with WebGL2 and keeps its adaptive resolution/fallback. Both smooth
+  renderers share `smoothSurface` for the cached parallax horizon.
+- Smooth classes are set in `cell.js`: 4 crisp built piece, 6 wall tile, 7/8 belt, 9–12 screw,
+  13/14 filter, 19–22 diagonal screw, 23–30 slope belt/filter, 31/32 slope sidewall and cleats.
+  `diagonal-screws.js` and `slope-belts.js` share the simulation join geometry with pixel colours,
+  smooth tags, ghosts and previews.
 
 ## The browser shell
 
@@ -421,29 +423,48 @@ Shared CPU colour rules, used by legacy and the rare-cell WebGPU override.
 Loose bars, flying grains and queued blueprint blocks in the grain buffer.
 - functions: worldObjects
 
-### src/render/gpu/colour.js (82 lines)
+### src/render/gpu/colour.js (101 lines)
 The plain-cell rules mirror cellPixel; built geometry stays in the CPU override.
-- variables: GPU_COLOUR
+- functions: gpuColourShader
+- variables: GPU_COLOUR, GPU_COLOUR_SMOOTH
 
 ### src/render/gpu/shaders.js (42 lines)
 Half-resolution separable Gaussian bloom, then nearest colour + linear light.
 - variables: GPU_BLUR, GPU_DRAW
 
-### src/render/gpu/device.js (43 lines)
+### src/render/gpu/light.js (143 lines)
+GPU sources retain the CPU roof and transmission rules; directional relaxation follows only left/up, then right/down, so it cannot introduce extra light paths.
+- functions: gpuLightGrid, gpuLightBindings, gpuLight
+- variables: GPU_LIGHT, GPU_LIGHT_SPREAD, GPU_LIGHT_PACK
+
+### src/render/gpu/smooth-slope.js (107 lines)
+Smooth slope geometry and cleats, ported from slope-belt-shader.js.
+- variables: GPU_SMOOTH_SLOPE_HELPERS, GPU_SMOOTH_SLOPE_DRAW
+
+### src/render/gpu/smooth-shader.js (368 lines)
+Literal port of the legacy smooth shader; dpdy is reversed for WebGPU's top-left origin.
+- variables: GPU_SMOOTH
+
+### src/render/gpu/smooth.js (47 lines)
+Field mip levels preserve the legacy trilinear AO sample at level 2.2.
+- functions: gpuSmoothInit, gpuSmoothTextures, gpuSmoothBindings, gpuSmoothUpload, gpuSmoothMips
+- variables: GPU_FIELD_MIP
+
+### src/render/gpu/device.js (48 lines)
 Async startup shares the title screen's adapter; failures leave a usable legacy renderer.
 - functions: gpuInit, gpuFail, worldPixelContext
 - variables: GPU
 
-### src/render/gpu/textures.js (84 lines)
+### src/render/gpu/textures.js (104 lines)
 Whole-world allocation, direct strided uploads; no cell packing or dirty tracking.
 - functions: gpuTexture, gpuWrite, gpuBuffer, gpuWorldTextures, gpuViewTextures, gpuBindings, gpuUpload
 
-### src/render/gpu/draw.js (69 lines)
+### src/render/gpu/draw.js (83 lines)
 Only exceptional blocks and their geometric neighbours run the CPU colour rules.
-- functions: gpuOverrides, gpuDraw, gpuPresent
+- functions: gpuOverrideHas, gpuOverrides, gpuDraw, gpuPresent
 - variables: GPU_PATTERN
 
-### src/render/world.js (83 lines)
+### src/render/world.js (85 lines)
 Scale colour without changing the smooth-look class.
 - functions: fogColor, drawWorld
 
@@ -456,9 +477,9 @@ Other overlays: floating text, rain and snow, flow arrows, filter tags, molds, t
 - functions: spawnFloater, drawRain, drawFlowArrows, drawJoinMark, drawJoinMarks, flowLoaded, drawFlowRates, pingImage, drawPing, drawHoveredVein, drawPingRecharge, drawRunOutline, drawFilterTags, drawFloaters, moldLabel, drawMolds, drawBot, spawnDebris, spawnSpark, spawnSuck, drawParticles
 - variables: PING_SPEED, PING_HOLD, PING_FADE, PING_BANDS, PING_GLINTS, GLINT_MS
 
-### src/render/smooth.js (103 lines)
+### src/render/smooth.js (108 lines)
 Smooth look: WebGL setup, texture uploads, drawing and adaptive resolution.
-- functions: smCompile, smoothAvailable, smUpload, smoothDraw, smoothPace
+- functions: smCompile, smoothAvailable, smUpload, smoothSurface, smoothDraw, smoothPace
 - variables: SM, smBytes
 
 ### src/render/slope-belt-shader.js (111 lines)
@@ -604,6 +625,7 @@ The frame loop and boot.
 
 - `test/firefox_smoke.py`: An isolated Firefox title screen and three seconds of a new pixel world.
 - `test/furnace_play.js`: Fires a starter furnace design repeatedly the way a player would, without cleaning it out between firings, to see whether it clogs. Usage: node test/furnace_play.js [design] [ore] [coal] [firings] [roll] [mode] mode: 'mix' (even ore/coal mix blown in), 'layers' (ore first, coal on top), 'vent' (blown in from above through the vent)
+- `test/gpu-readback.js`: Read the four actual smooth textures without adding COPY_SRC usage to game resources.
 - `test/harness.js`: Loads the simulation (the src/sim/ files listed in src/manifest.json, joined in order into one script, the same way build.py joins them) so scenarios can drive it headlessly in Node. Errors point at the real file and line (sim/pipes.js:42), not at the joined script.
 - `test/play.py`: Play-test in a real browser: mouse, keyboard and a mocked gamepad drive the game; save and continue. Screenshots go to shots/play/ (or the folder given as the first argument). Pixel look unless the second argument is 'smooth'.
 - `test/run_all.py`: Run the test suites and report.
@@ -637,7 +659,7 @@ The frame loop and boot.
 - `test/t_slope_belt_art.js`: Drawing geometry against real parts, without changing transport.
 - `test/t_slope_belts.js`: Sloped belts: transport, joins, sorting, persistence and the actual drag controls.
 - `test/t_slopes.js`: Sloped walls: flow, sealed joints, building, refunds and persistence.
-- `test/t_smooth.py`: The smooth look: the WebGL pass compiles, draws, falls back cleanly, and the thermal lens still works.
+- `test/t_smooth.py`: The smooth look: the WebGPU pass compiles, draws, falls back cleanly, and the thermal lens still works.
 - `test/t_stone.js`: Uses for waste rock and sand: glass bricks from a fan furnace, the Saw, the silo and builders' orders.
 - `test/t_water.js`: Water conservation, settling, ledge pouring, boundaries, pipes and phase changes. Tall worlds keep geothermal heat out of the ponds.
 - `test/t_weather.js`: Weather: no more springs; steam rises into clouds and rains back; showers; snow when it is freezing.
@@ -679,4 +701,5 @@ The frame loop and boot.
 - `tools/scenes/t35-joins.js`: Run after clear-sky.js. Paused, empty pipes make casing, flights and mouths easy to inspect.
 - `tools/scenes/t36-funnel.js`: A twelve-block funnel, one block open at the bottom, draining 400 grains of sand.
 - `tools/scenes/t40-sloped-belt.js`: Self-contained: a flat feed climbs, crosses the top, and runs downhill; a filter sorts below.
+- `tools/scenes/t88-light.js`: Frozen light comparisons, also used by the WebGPU browser regression suite.
 - `tools/scenes/t88-renderer.js`: Paused comparison scene: ordinary grains, fog, heat, trees and exceptional geometry.

docs/TESTING.md

@@ -67,7 +67,7 @@ diff -r /tmp/before /tmp/after
   gamepad and world sizes. Screenshots go to
   `shots/play/`. Pixel look by default; `python3 test/play.py shots/play smooth` uses the smooth
   look.
-- `test/t_smooth.py` checks that the WebGL2 smooth look compiles and draws, that it falls back
+- `test/t_smooth.py` checks that the WebGPU smooth look compiles and draws, compares legacy at half resolution, and falls back
   to pixels cleanly, and that the thermal lens works.
 - Run them one at a time. Two headless browsers at once can starve each other, and
   `t_smooth.py`'s half-resolution check then fails. `run_all.py` runs them one after the other.
@@ -110,8 +110,8 @@ python3 tools/shot.py shots/silo.png --setup tools/scenes/clear-sky.js \
   `tools/scenes/clear-sky.js` clears a flat work area and sets `FLOOR`, `BX`, `BY` and `CAM`. Set
   `CAM = [x, y]` to frame the shot.
 - The HUD and the cursor previews are hidden unless you pass `--hud` or `--previews`.
-- `--renderer webgpu|legacy` selects the pixel renderer (default WebGPU). WebGPU screenshots
+- `--renderer webgpu|legacy` selects the renderer for either look (default WebGPU). WebGPU screenshots
-  fail clearly when no device is available; Smooth still uses WebGL2.
+  fail clearly when no device is available; Legacy smooth uses WebGL2; WebGPU smooth uses WGSL.
 - `test/t_gpu.py` compares the rendered composite against Legacy for ordinary cells, built pieces,
   ghosts, fog and all thermal palettes, and checks world replacement, resize, timing and device loss.
   World-pixel tests use `worldPixelContext()` to read both stacked canvases. The helper
@@ -233,3 +233,15 @@ CPU. The GPU canvas draws the world; the transparent game canvas holds the exist
 The comparison fixture is `tools/scenes/t88-renderer.js`. Capture with `--freeze --zoom 12`
 (and 6 or 2), `--viewport 1920x1080`, and each `--renderer`. It includes terrain, water, fire,
 magma, fog, a tree, flat/slope belts, screws, pipe, mold, bar and blueprint ghost.
+
+WebGPU Phase 2c: both looks use WebGPU. `test/t_gpu.py` also reads all four smooth input
+textures back and compares them against the legacy CPU buffers at zoom 12/6/2 with fog on/off:
+classes, sky flags, field bits and geometry tags must match exactly; RGB can differ by one.
+The smooth shader retains the legacy nine-tap glow filter and samples field AO at mip level 2.2;
+three render passes generate field levels 1-3, including non-power-of-two view sizes.
+Only smooth allocates/uploads the back/tree/field overrides and a separate presence mask.
+`tools/perf.py` refuses a WebGPU fallback in either look and records WebGPU scale as 1.
+The screenshot tool invalidates the shared parallax horizon after scene setup changes `surfY`,
+so renderer startup timing cannot change the hills in a comparison. Mip pass references:
+[Chrome's porting guide](https://developer.chrome.com/docs/web-platform/webgpu/from-webgl-to-webgpu#mipmap_generation)
+and [WebGPU Fundamentals](https://webgpufundamentals.org/webgpu/lessons/webgpu-importing-textures.html#generating-mips-on-the-gpu).

docs/agent-log.md

@@ -4,6 +4,20 @@ What agents have done, newest first. Add an entry for every piece of work: the d
 note or request it was for, what changed, the files touched, the tests run, and anything left
 open or not verified. Keep each entry short; the git history has the details.
 
+## 2026-10-27: t88 phase 2c, smooth look on WebGPU (Codex + Claude Code)
+
+- WGSL port of the smooth look (`src/render/gpu/smooth.js`, `smooth-shader.js`, `smooth-slope.js`):
+  the GPU colour pass writes all four smooth inputs for plain cells; the CPU rare-cell overrides
+  (`cell.js`, now also mold walls) fill them for built pieces and objects (12.73/3.41/0.38% of
+  factory cells at zoom 12/6/2). Field AO via three GPU mip passes; nine-tap glow kept; no
+  WebGPU `smoothPace` (max GPU time 1.24 ms). Legacy still selectable.
+- Checked by Codex: 21 screenshot pairs legacy vs WebGPU, mean RGB error 0.01-0.22/255, small
+  counts off by more than 8 (thin AO and edge rounding), no missing shapes or open joins; input
+  readbacks equal legacy. Smooth factory drawWorld zoom 2 28.2 -> 10.0 ms; real loop zoom 2 CPU
+  p50/p95 34.7/39.0 -> 14.5/15.7 ms. Claude did not inspect the images (token budget).
+- Open for 2d: adapters with fewer than five storage textures per stage fall back to legacy.
+- Tests: `python3 test/run_all.py --browser` passed.
+
 ## 2026-10-27: t88 phase 2b, lighting on the GPU (Codex + Claude Code)
 
 - `src/render/gpu/light.js`: light sources per light cell (transmittance, sky from `roofY`, heat,

src/manifest.json

@@ -53,6 +53,9 @@
     "render/gpu/colour.js",
     "render/gpu/shaders.js",
     "render/gpu/light.js",
+    "render/gpu/smooth-slope.js",
+    "render/gpu/smooth-shader.js",
+    "render/gpu/smooth.js",
     "render/gpu/device.js",
     "render/gpu/textures.js",
     "render/gpu/draw.js",

src/render/cell.js

@@ -2,7 +2,7 @@
 function cellPixel(x, y) {
   const { bw, bh, px, vx0, vy0, smoothOn: sm, rain: rk } = R;
   const bx = x - vx0, by = y - vy0, row = by * bw;
-  if (R.gpuOn && !px[row + bx]) GPU.overrideCount++;
+  if (R.gpuOn) { if (!gpuOverrideHas(row + bx)) GPU.overrideCount++; if (sm) GPU.overrideMask[row + bx] = 255; }
   const bpx = R.backPx, tpx = R.treePx, fpx = R.fieldPx, mmask = R.moldMask;
   const OUT = sm ? 0x04000000 : 0xff000000;
   if (x < 0 || y < 0 || x >= W || y >= H) {

src/render/gpu/colour.js

@@ -1,5 +1,5 @@
 // The plain-cell rules mirror cellPixel; built geometry stays in the CPU override.
-const GPU_COLOUR = `
+function gpuColourShader(smooth) { return `
 struct Params { view: vec4f, world: vec4f, style: vec4f, effects: vec4f, camera: vec4f, light: vec4f, screen: vec4f, extra: vec4f }
 @group(0) @binding(0) var<uniform> u: Params;
 @group(0) @binding(1) var<storage, read> lut: array<vec4f>;
@@ -9,6 +9,14 @@ ${['mat','flags','life','part','bgm','digp','noise','fog'].map((n,i)=>`@group(0)
 @group(0) @binding(12) var overrideT: texture_2d<f32>;
 @group(0) @binding(13) var overrideGlow: texture_2d<f32>;
 @group(0) @binding(14) var outputT: texture_storage_2d<rgba8unorm, write>;
+${smooth ? `@group(0) @binding(16) var backOut: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(17) var treeOut: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(18) var fieldOut: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(19) var backOver: texture_2d<f32>;
+@group(0) @binding(20) var treeOver: texture_2d<f32>;
+@group(0) @binding(21) var fieldOver: texture_2d<f32>;
+@group(0) @binding(22) var maskOver: texture_2d<u32>;
+` : ""}
 @group(0) @binding(15) var glowT: texture_storage_2d<rgba8unorm, write>;
 fn phase(p: vec2i) -> i32 { return i32(lut[textureLoad(matT,p,0).r*3u+1u].x); }
 fn powder(p: vec2i) -> bool {
@@ -20,6 +28,7 @@ fn colour(@builtin(global_invocation_id) id: vec3u) {
   let q=vec2i(id.xy); if (any(q>=vec2i(u.view.zw))) { return; }
   let p=q+vec2i(u.view.xy); let x=p.x; let y=p.y;
   var rgb=vec3f(0); var bloom=vec3f(0);
+  ${smooth ? "var cl=4.; var backValue=vec4f(0); var treeValue=vec4f(0); var fieldValue=vec4f(1,0,0,0);" : ""}
   if (all(p>=vec2i(0)) && all(p<vec2i(u.world.xy))) {
     let m=textureLoad(matT,p,0).r; let nz=f32(textureLoad(noiseT,p,0).r)-128.;
     let T=textureLoad(tempT,p,0).r; let mc=lut[m*3u]; let props=lut[m*3u+1u];
@@ -32,7 +41,7 @@ fn colour(@builtin(global_invocation_id) id: vec3u) {
     if (fv>0.) {
       var shade=1.;
       let built=props.x==${PH_SOLID}. && (textureLoad(flagsT,p,0).r & ${F_NAT}u)==0u;
-      if (m!=${EMPTY}u && !built) {
+      if (${smooth ? "false" : `m!=${EMPTY}u && !built`}) {
         if (props.x==${PH_POWDER}.) { shade=1.1+nz*0.00065; }
         else if (props.x==${PH_SOLID}. && (powder(p+vec2i(-1,0)) || powder(p+vec2i(1,0)) || powder(p+vec2i(0,-1)) || powder(p+vec2i(0,1)))) { shade=0.72; }
       }
@@ -43,7 +52,7 @@ fn colour(@builtin(global_invocation_id) id: vec3u) {
         let tb=textureLoad(bgmT,p,0).r; let dp=f32(textureLoad(digpT,p,0).r);
         let column=land[x]; let sky=land[i32(u.world.x)+y].rgb;
         var back: vec3f;
-        if (tb!=0u) { back=(lut[tb*3u+2u].rgb+vec3f(nz*lut[tb*3u].a/128.))*(1.+dp/700.); }
+        if (${smooth ? "false" : "tb!=0u"}) { back=(lut[tb*3u+2u].rgb+vec3f(nz*lut[tb*3u].a/128.))*(1.+dp/700.); }
         else if (f32(y)<column.x) {
           back=sky;
           if (u.style.y>0.01) { let gy=(sky.r+sky.g+sky.b)/3.*0.82; back+= (vec3f(gy)-sky)*u.style.y*vec3f(0.7,0.7,0.65); }
@@ -52,13 +61,18 @@ fn colour(@builtin(global_invocation_id) id: vec3u) {
           if (f32(y)<column.y) { back=vec3f(${BG.dirt.join(',')}); }
           else if (f32(y)<column.z) { back=vec3f(${BG.stone.join(',')}); }
           else if (f32(y)<column.w) { back=vec3f(${BG.granite.join(',')}); }
-          back+=vec3f(nz*0.06);
+          ${smooth ? "" : "back+=vec3f(nz*0.06);"}
         }
+        ${smooth ? `backValue=vec4f(quant(back),select(0.,1.,f32(y)<column.x));
+        if (tb!=0u) { treeValue=vec4f(quant((lut[tb*3u+2u].rgb+vec3f(nz*lut[tb*3u].a/256.))*(1.+dp/700.)),select(128./255.,1.,tb==${LEAVES}u)); }
+        cl=0.; if (m!=${EMPTY}u) { cl=5.; if(props.x==${PH_SOLID}.) { cl=1.; } else if(props.x==${PH_POWDER}.) { cl=2.; } else if(props.x==${PH_LIQUID}.) { cl=3.; } }
+        fieldValue=vec4f(select(0.,1.,cl==1.),select(0.,1.,cl==3.),select(0.,1.,cl==5.),select(0.,1.,tb!=0u));
+        ` : ""}
         rgb=back;
         if (m!=${EMPTY}u) {
-          rgb=(mc.rgb+vec3f(nz*mc.a*u.style.z/128.))*shade;
+          rgb=(mc.rgb+vec3f(nz*mc.a*${smooth ? `select(u.style.z*0.35,0.,props.x==${PH_SOLID}.)` : "u.style.z"}/128.))*shade;
           if (m==${FIRE}u) { let f=f32(textureLoad(lifeT,p,0).r & 15u)*6.; rgb=vec3f(255.,120.+f,30.+f*0.5); }
-          else if (m==${WATER}u && ((x*3+y*5+(i32(u.style.w)>>2)) & 15)==0) { rgb+=vec3f(14.); }
+          else if (${smooth ? "false && " : ""}m==${WATER}u && ((x*3+y*5+(i32(u.style.w)>>2)) & 15)==0) { rgb+=vec3f(14.); }
           rgb=back+(rgb-back)*props.y;
           if (dp>0.) { let lift=dp/255.*(30.+sin(u.style.w*1.9+f32(x)*1.3+f32(y)*0.7)*6.); rgb+=lift*vec3f(1,1,1.08); }
         }
@@ -72,11 +86,16 @@ fn colour(@builtin(global_invocation_id) id: vec3u) {
         }
       }
       rgb=quant(rgb*fv/255.);
+      ${smooth ? "backValue=vec4f(quant(backValue.rgb*fv),backValue.a); treeValue=vec4f(quant(treeValue.rgb*fv),treeValue.a);" : ""}
     }
   }
   // Flying grains replace colour, while the glow behind them remains.
   let over=textureLoad(overrideT,q,0);
-  if (over.a>0.) { rgb=over.rgb; bloom=max(bloom,textureLoad(overrideGlow,q/2,0).rgb); }
+  if (${smooth ? "textureLoad(maskOver,q,0).r>0u" : "over.a>0."}) { rgb=over.rgb; ${smooth ? "cl=over.a*255.; backValue=textureLoad(backOver,q,0); treeValue=textureLoad(treeOver,q,0); fieldValue=textureLoad(fieldOver,q,0);" : ""} bloom=max(bloom,textureLoad(overrideGlow,q/2,0).rgb); }
-  textureStore(outputT,q,vec4f(rgb,1));
+  ${smooth ? "textureStore(backOut,q,backValue); textureStore(treeOut,q,treeValue); textureStore(fieldOut,q,fieldValue);" : ""}
+  textureStore(outputT,q,vec4f(rgb,${smooth ? "cl/255." : "1"}));
   if (((q.x | q.y) & 1)==0) { textureStore(glowT,q/2,vec4f(bloom,1)); }
 }`;
+}
+const GPU_COLOUR = gpuColourShader(false);
+const GPU_COLOUR_SMOOTH = gpuColourShader(true);

src/render/gpu/device.js

@@ -3,7 +3,7 @@ const GPU = { ok: false, canvas: document.getElementById('gpu'), error: '', gpuM
 async function gpuInit(adapter) {
   try {
     const features = adapter.features.has('timestamp-query') ? ['timestamp-query'] : [];
-    const d = GPU.device = await adapter.requestDevice({ requiredFeatures: features });
+    const d = GPU.device = await adapter.requestDevice({ requiredFeatures: features, requiredLimits: { maxStorageTexturesPerShaderStage: 5 } });
     d.addEventListener('uncapturederror', e => gpuFail(e.error.message));
     d.lost.then(info => gpuFail(info.message || 'Graphics device lost'));
     GPU.context = GPU.canvas.getContext('webgpu');
@@ -11,6 +11,8 @@ async function gpuInit(adapter) {
     GPU.context.configure({ device: d, format: GPU.format, alphaMode: 'opaque' });
     d.pushErrorScope('validation');
     GPU.colour = await d.createComputePipelineAsync({ layout: 'auto', compute: { module: d.createShaderModule({ code: GPU_COLOUR }), entryPoint: 'colour' } });
+    GPU.colourSmooth = await d.createComputePipelineAsync({ layout: 'auto', compute: { module: d.createShaderModule({ code: GPU_COLOUR_SMOOTH }), entryPoint: 'colour' } });
+    gpuSmoothInit();
     for (const [name, code, entryPoint] of [['lightRoof',GPU_LIGHT,'roof'],['lightSources',GPU_LIGHT,'sources'],['lightSpread',GPU_LIGHT_SPREAD,'main'],['lightPack',GPU_LIGHT_PACK,'main']]) {
       GPU[name] = await d.createComputePipelineAsync({layout:'auto',compute:{module:d.createShaderModule({code}),entryPoint}});
     }
@@ -19,7 +21,7 @@ async function gpuInit(adapter) {
     const error = await d.popErrorScope(); if (error) throw error;
     GPU.uniform = d.createBuffer({ size: 128, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST });
     GPU.params = new Float32Array(32);
-    GPU.sampler = d.createSampler({ magFilter: 'linear', minFilter: 'linear' });
+    GPU.sampler = d.createSampler({ magFilter: 'linear', minFilter: 'linear', mipmapFilter: 'linear' });
     if (features.length) {
       GPU.queries = d.createQuerySet({ type: 'timestamp', count: 4 });
       GPU.resolve = d.createBuffer({ size: 256, usage: GPUBufferUsage.QUERY_RESOLVE | GPUBufferUsage.COPY_SRC });

src/render/gpu/draw.js

@@ -1,8 +1,13 @@
 // Only exceptional blocks and their geometric neighbours run the CPU colour rules.
 const GPU_PATTERN = new Uint8Array(NMAT);
-for (const m of [BELTL, BELTR, PIPE, SHAFT, GATE, CUTST, GRINDER, SAWB, PRESSB, RACKB, WOOD, LWOOD, HULL, SPRING]) GPU_PATTERN[m] = 1;
+for (const m of [BELTL, BELTR, PIPE, SHAFT, GATE, CUTST, GRINDER, SAWB, PRESSB, RACKB, WOOD, LWOOD, HULL, SPRING, DOCKW]) GPU_PATTERN[m] = 1;
+function gpuOverrideHas(k) { return R.smoothOn ? GPU.overrideMask[k] : R.px[k]; }
 function gpuOverrides() {
   R.px.fill(0);
+  if (R.smoothOn) {
+    if (!GPU.overrideMask || GPU.overrideMask.length !== R.px.length) GPU.overrideMask = new Uint8Array(R.px.length);
+    GPU.overrideMask.fill(0);
+  }
   if (!GPU.blocks || GPU.blocks.length !== NB) GPU.blocks = new Uint8Array(NB);
   const marked = GPU.blocks; marked.fill(0);
   const bx0 = Math.max(0, (R.vx0 >> 2) - 3), by0 = Math.max(0, (R.vy0 >> 2) - 3);
@@ -22,6 +27,10 @@ function gpuOverrides() {
         for (let xx = Math.max(0, bx - 2); xx <= Math.min(BW - 1, bx + 2); xx++) marked[yy * BW + xx] = 1;
     }
   }
+  if (R.smoothOn) for (const mo of molds) {
+    if (mo.x + 40 < R.vx0 || mo.x - 40 > R.vx0 + R.bw || mo.y + 40 < R.vy0 || mo.y - 40 > R.vy0 + R.bh) continue;
+    for (const i of moldCellsAt(mo.kind, mo.x, mo.y).walls) marked[(((i / W) | 0) >> 2) * BW + ((i % W) >> 2)] = 1;
+  }
   GPU.overrideCount = 0;
   for (let by = by0; by <= by1; by++) for (let bx = bx0; bx <= bx1; bx++) {
     if (!marked[by * BW + bx]) continue;
@@ -39,14 +48,15 @@ function gpuDraw(lighting, glowOn, lightUpdate) {
   const blit = performance.now(), d = GPU.device, encoder = d.createCommandEncoder();
   const timed = GPU.queries && !GPU.timingBusy && (PERF_MARKS || SET.showPerf);
   const compute = encoder.beginComputePass(timed ? { timestampWrites: { querySet: GPU.queries, beginningOfPassWriteIndex: 0 } } : {});
-  compute.setPipeline(GPU.colour); compute.setBindGroup(0, GPU.bindColour);
+  compute.setPipeline(R.smoothOn ? GPU.colourSmooth : GPU.colour); compute.setBindGroup(0, R.smoothOn ? GPU.bindColourSmooth : GPU.bindColour);
   compute.dispatchWorkgroups(Math.ceil(R.bw / 8), Math.ceil(R.bh / 8)); compute.end();
-  if (glowOn) for (let i = 0; i < 2; i++) {
+  if (glowOn && !R.smoothOn) for (let i = 0; i < 2; i++) {
     const pass = encoder.beginComputePass(); pass.setPipeline(GPU.blur); pass.setBindGroup(0, GPU.bindBlur[i]);
     pass.dispatchWorkgroups(Math.ceil(R.glowC.width / 8), Math.ceil(R.glowC.height / 8)); pass.end();
   }
   const lightTimed = lighting && (lightUpdate || GPU.lightFrame === undefined);
   if (lightTimed) { gpuLight(encoder, timed); GPU.lightFrame = R.frame; }
+  if (R.smoothOn) gpuSmoothMips(encoder);
   gpuPresent(encoder, timed);
   if (timed) { encoder.resolveQuerySet(GPU.queries, 0, lightTimed ? 4 : 2, GPU.resolve, 0); encoder.copyBufferToBuffer(GPU.resolve, 0, GPU.readback, 0, lightTimed ? 32 : 16); }
   d.queue.submit([encoder.finish()]);
@@ -69,5 +79,5 @@ function gpuDraw(lighting, glowOn, lightUpdate) {
 function gpuPresent(encoder, timed = false) {
   const pass = encoder.beginRenderPass({ colorAttachments: [{ view: GPU.context.getCurrentTexture().createView(), loadOp: 'clear', storeOp: 'store', clearValue: [0, 0, 0, 1] }],
     ...(timed ? { timestampWrites: { querySet: GPU.queries, endOfPassWriteIndex: 1 } } : {}) });
-  pass.setPipeline(GPU.draw); pass.setBindGroup(0, GPU.bindDraw); pass.draw(3); pass.end();
+  pass.setPipeline(R.smoothOn ? GPU.smooth : GPU.draw); pass.setBindGroup(0, R.smoothOn ? GPU.bindSmooth : GPU.bindDraw); pass.draw(3); pass.end();
 }

src/render/gpu/smooth-shader.js

@@ -0,0 +1,368 @@
+// Literal port of the legacy smooth shader; dpdy is reversed for WebGPU's top-left origin.
+const GPU_SMOOTH = `
+struct Params { view: vec4f, world: vec4f, style: vec4f, effects: vec4f, camera: vec4f, light: vec4f, screen: vec4f, extra: vec4f }
+@group(0) @binding(0) var<uniform> u: Params;
+@group(0) @binding(1) var tCol: texture_2d<f32>;
+@group(0) @binding(2) var tBack: texture_2d<f32>;
+@group(0) @binding(3) var tTree: texture_2d<f32>;
+@group(0) @binding(4) var tField: texture_2d<f32>;
+@group(0) @binding(5) var tLight: texture_2d<f32>;
+@group(0) @binding(6) var tGlow: texture_2d<f32>;
+@group(0) @binding(7) var linear: sampler;
+@group(0) @binding(8) var<uniform> su: vec4f;
+fn sampleLevel(t: texture_2d<f32>, uv: vec2f, level: f32) -> vec4f { return textureSampleLevel(t,linear,uv,level); }
+fn modulo(x: f32,y: f32) -> f32 { return x-y*floor(x/y); }
+fn h21(inputP: vec2f) -> f32 { var p=inputP; p = fract(p * vec2f(123.34, 456.21)); p += dot(p, p + 45.32); return fract(p.x * p.y); }
+fn vn(p: vec2f) -> f32 {
+  var i: vec2f = floor(p); var f: vec2f = fract(p); f = f * f * (3.0 - 2.0 * f);
+  return mix(mix(h21(i), h21(i + vec2f(1, 0)), f.x), mix(h21(i + vec2f(0, 1)), h21(i + vec2f(1, 1)), f.x), f.y);
+}
+fn fbm(inputP: vec2f) -> f32 { var p=inputP; var a: f32 = 0.5; var s: f32 = 0.0; for (var k: i32 = 0; k < 4; k++) { s += a * vn(p); p = p * 2.03 + 17.1; a *= 0.5; } return s / 0.9375; }
+
+fn lin(t: texture_2d<f32>, q: vec2f) -> vec4f { return sampleLevel(t, q / (u.view.zw), 0.0); }
+
+fn bsp(t: texture_2d<f32>, q: vec2f) -> vec4f {
+  var tc: vec2f = q - 0.5; var i: vec2f = floor(tc); var f: vec2f = tc - i; var f2: vec2f = f * f; var f3: vec2f = f2 * f;
+  var w0: vec2f = (1.0 - 3.0 * f + 3.0 * f2 - f3) / 6.0; var w1: vec2f = (4.0 - 6.0 * f2 + 3.0 * f3) / 6.0;
+  var w2: vec2f = (1.0 + 3.0 * f + 3.0 * f2 - 3.0 * f3) / 6.0; var w3: vec2f = f3 / 6.0;
+  var g0: vec2f = w0 + w1; var g1: vec2f = w2 + w3;
+  var h0: vec2f = (i - 0.5 + w1 / g0) / (u.view.zw); var h1: vec2f = (i + 1.5 + w3 / g1) / (u.view.zw);
+  return g0.y * (g0.x * sampleLevel(t, h0, 0.0) + g1.x * sampleLevel(t, vec2f(h1.x, h0.y), 0.0))
+       + g1.y * (g0.x * sampleLevel(t, vec2f(h0.x, h1.y), 0.0) + g1.x * sampleLevel(t, h1, 0.0));
+}
+fn cellAt(t: texture_2d<f32>, c: vec2i) -> vec4f { return textureLoad(t, clamp(c, vec2i(0), vec2i((u.view.zw)) - 1), 0); }
+fn cls(c: vec4f) -> i32 { return i32(c.a * 255.0 + 0.5); }
+fn isT(n: i32) -> f32 { return select( 0.0, 1.0 ,(n == 1 || n == 4 || n == 6 || (n >= 7 && n <= 12)) ); }
+fn isM(n: i32) -> f32 { return select( isT(n), 1.0 ,n == 2 ); }   
+fn isF(n: i32) -> f32 { return select( 0.0, 1.0 ,(n == 13 || n == 14) ); }
+fn isBelt(n: i32) -> f32 { return select( 0.0, 1.0 ,(n == 7 || n == 8) ); }
+fn isN(n: i32, k: i32) -> f32 { return select( 0.0, 1.0 ,n == k ); }
+
+fn beltSD(lb: vec2f, endL: bool, endR: bool) -> f32 {
+  var x0: f32 = select( -2.0, 0.15 , endL ); var x1: f32 = select( 6.0, 3.85 , endR ); var r: f32 = 0.75;
+  var c: vec2f = vec2f((x0 + x1) * 0.5, 2.0); var hs: vec2f = vec2f((x1 - x0) * 0.5, 1.85);
+  var q: vec2f = abs(lb - c) - hs + r;
+  return length(max(q, vec2f(0.0))) + min(max(q.x, q.y), 0.0) - r;
+}
+
+fn beltLineSD(c: vec2i, lb: vec2f) -> f32 {
+  var caps: i32 = i32(cellAt(tTree, c).a * 255.0 + 0.5);
+  return beltSD(lb, (caps & 1) != 0, (caps & 2) != 0);
+}
+
+fn bandRidge(w: vec2f, lb: vec2f, dir: f32) -> f32 {
+  var run: f32 = select( -1.0, 1.0 , lb.y < 2.0 );
+  return step(0.72, fract((w.x - dir * run * (su.x) * 30.0) * 0.5));
+}
+fn bandColor(lb: vec2f, rg: f32) -> vec3f { return mix(vec3f(0.2, 0.21, 0.24), vec3f(0.34, 0.35, 0.39), rg) * (select( 0.72, 1.0 ,lb.y < 2.0 )); }
+
+${GPU_SMOOTH_SLOPE_HELPERS}
+@fragment fn fs(@builtin(position) position: vec4f) -> @location(0) vec4f {
+  let uView=vec4f(u.extra.xy-u.view.xy,u.camera.z,u.screen.y);
+  let uLightMap=vec4f(u.light.xy*u.light.z-u.view.xy,u.light.z,u.light.w);
+  var zp: f32 = uView.z; var px: f32 = 1.0 / zp;
+  var sp: vec2f = vec2f(position.x, position.y);
+  var q: vec2f = sp * px + uView.xy;          
+  var w: vec2f = q + (u.view.xy);                 
+  var cam: vec2f = (u.view.xy) + uView.xy;        
+  var scr: vec2f = q - uView.xy;              
+  var shc: f32 = uView.w * px;             
+  var c: vec2i = vec2i(floor(q));
+  var qq: vec2f = q - 0.5; var c0: vec2i = vec2i(floor(qq)); var t: vec2f = qq - floor(qq);
+  var k0: vec4f = cellAt(tCol, c0); var k1: vec4f = cellAt(tCol, c0 + vec2i(1, 0)); var k2: vec4f = cellAt(tCol, c0 + vec2i(0, 1)); var k3: vec4f = cellAt(tCol, c0 + vec2i(1, 1));
+  var n0: i32 = cls(k0); var n1: i32 = cls(k1); var n2: i32 = cls(k2); var n3: i32 = cls(k3);
+  var own: vec4f = cellAt(tCol, c); var nc: i32 = cls(own);
+  var s: vec2f = smoothstep(vec2f(0.1), vec2f(0.9), t);
+  var wq: vec4f = vec4f((1.0 - s.x) * (1.0 - s.y), s.x * (1.0 - s.y), (1.0 - s.x) * s.y, s.x * s.y);
+
+  // Fields and screen derivatives before branching.
+  var Fb: vec4f = bsp(tField, q);
+  var F: vec4f = mix(lin(tField, q), Fb, 0.62);
+  
+  var wl: vec4f = vec4f((1.0 - t.x) * (1.0 - t.y), t.x * (1.0 - t.y), (1.0 - t.x) * t.y, t.x * t.y);
+  var pw: f32 = dot(wl, vec4f(isN(n0, 2), isN(n1, 2), isN(n2, 2), isN(n3, 2)));
+  var M: f32 = clamp(F.r + pw, 0.0, 1.0);
+  var lf: f32 = clamp(mix(F.g, Fb.g, 0.6) + F.r + pw, 0.0, 1.0);
+  var fwT: f32 = max(fwidth(F.r), 1e-4); var fwM: f32 = max(fwidth(M), 1e-4); var fwA: f32 = max(fwidth(F.a), 1e-4); var fwL: f32 = max(fwidth(lf), 1e-4);
+  var gT: vec2f = vec2f(dpdx(mix(Fb.r, M, 0.5)), -dpdy(mix(Fb.r, M, 0.5))); var gL: vec2f = vec2f(dpdx(F.g), -dpdy(F.g));
+
+  // Sky, clouds, parallax hills and back-wall AO.
+  var back: vec4f = lin(tBack, q);
+  var C: vec3f = back.rgb;
+  var sky: f32 = back.a;
+  if (sky > 0.01) {
+    var S: vec3f = C;
+    var P0: vec2f = scr + cam * 0.08;
+    
+    var cl: f32 = smoothstep(0.56 - 0.3 * (u.style.y), 0.78 - 0.2 * (u.style.y), fbm(vec2f(P0.x * 0.016 + (su.x) * 0.006, P0.y * 0.05)));
+    var hi: f32 = clamp(((su.y) - 40.0 - w.y) / 50.0, 0.0, 1.0);
+    S = mix(S, mix(vec3f(0.96, 0.95, 0.93), vec3f(0.5, 0.52, 0.56), (u.style.y)), cl * (0.6 + 0.3 * (u.style.y)) * hi);
+    var k1: f32 = 0.16; var P1: vec2f = scr + cam * k1;
+    var base1: f32 = shc * 0.5 + 8.0 + ((su.y) - shc * 0.5) * k1;
+    var top1: f32 = base1 - 36.0 * (0.25 + 0.9 * fbm(vec2f(P1.x * 0.009, 1.7)));
+    var m1: f32 = smoothstep(top1 - px * 1.5, top1 + px * 1.5, P1.y);
+    var far: vec3f = mix(vec3f(0.50, 0.60, 0.72), S, 0.35) * (0.94 + 0.08 * fbm(P1 * 0.05));
+    S = mix(S, far, m1);
+    var k2: f32 = 0.38; var P2: vec2f = scr + cam * k2;
+    var base2: f32 = shc * 0.5 + 14.0 + ((su.y) - shc * 0.5) * k2;
+    var top2: f32 = base2 - 20.0 * (0.2 + 0.9 * fbm(vec2f(P2.x * 0.016, 7.3)));
+    var m2: f32 = smoothstep(top2 - px * 1.5, top2 + px * 1.5, P2.y);
+    var near: vec3f = vec3f(0.36, 0.5, 0.47) * (0.92 + 0.14 * fbm(P2 * 0.08));
+    near = mix(near, S, 0.42 - 0.2 * smoothstep(top2, top2 + 30.0, P2.y));
+    S = mix(S, near, m2);
+    C = mix(C, S, sky);
+  }
+  
+  var notSky: f32 = 1.0 - sky;
+  C *= mix(1.0, 1.08 + 0.4 * fbm(w * 0.21), notSky);
+  var near: f32 = sampleLevel(tField, q / (u.view.zw), 2.2).r;
+  C *= 1.0 - 0.34 * smoothstep(0.12, 0.85, near) * (1.0 - 0.75 * sky);
+
+  // Trees behind terrain and built pieces.
+  if (F.a > 0.2) {
+    var a0: vec4f = cellAt(tTree, c0); var a1: vec4f = cellAt(tTree, c0 + vec2i(1, 0)); var a2: vec4f = cellAt(tTree, c0 + vec2i(0, 1)); var a3: vec4f = cellAt(tTree, c0 + vec2i(1, 1));
+    var m: vec4f = vec4f(step(0.4, a0.a), step(0.4, a1.a), step(0.4, a2.a), step(0.4, a3.a)) * wq;
+    var sm: f32 = m.x + m.y + m.z + m.w;
+    if (sm > 1e-4) {
+      var tc: vec3f = (a0.rgb * m.x + a1.rgb * m.y + a2.rgb * m.z + a3.rgb * m.w) / sm;
+      var leaf: f32 = (a0.a * m.x + a1.a * m.y + a2.a * m.z + a3.a * m.w) / sm;
+      var lv: f32 = smoothstep(0.6, 0.9, leaf);
+      tc *= mix(0.86 + 0.22 * vn(vec2f(w.x * 1.4, w.y * 0.25)), 0.72 + 0.5 * fbm(w * 0.42), lv);
+      var d: f32 = (F.a - 0.36) / fwA;
+      tc *= mix(0.55, 1.0, smoothstep(0.4, 1.6, d));
+      C = mix(C, tc, smoothstep(-0.7, 0.7, d));
+    }
+  }
+
+  // Smoke, steam and flame.
+  if (Fb.b > 0.02) {
+    var m: vec4f = vec4f(isN(n0, 5), isN(n1, 5), isN(n2, 5), isN(n3, 5)) * wq;
+    var sm: f32 = m.x + m.y + m.z + m.w;
+    var gc: vec3f = select( own.rgb, (k0.rgb * m.x + k1.rgb * m.y + k2.rgb * m.z + k3.rgb * m.w) / sm , sm > 1e-4 );
+    if (sm > 1e-4 || nc == 5) { C = mix(C, gc, smoothstep(0.03, 0.55, Fb.b) * 0.92); }
+  }
+
+  // Liquids with a glossy top edge.
+  var anyL: f32 = isN(n0, 3) + isN(n1, 3) + isN(n2, 3) + isN(n3, 3);
+  if (anyL > 0.0 && lf > 0.3) {
+    var m: vec4f = vec4f(isN(n0, 3), isN(n1, 3), isN(n2, 3), isN(n3, 3)) * (wq + 0.02);
+    var sm: f32 = m.x + m.y + m.z + m.w;
+    var lc: vec3f = (k0.rgb * m.x + k1.rgb * m.y + k2.rgb * m.z + k3.rgb * m.w) / sm;
+    lc *= 0.94 + 0.1 * vn(w * 0.55 + vec2f((su.x) * 0.7, (su.x) * 0.25));
+    var gl: f32 = length(gL);
+    var topSide: f32 = select( 0.0, clamp(gL.y / gl, 0.0, 1.0) , gl > 1e-5 );
+    var band: f32 = 1.0 - smoothstep(0.5, 0.75, F.g);
+    lc += vec3f(0.16, 0.18, 0.2) * topSide * band;
+    C = mix(C, lc, smoothstep(0.5 - fwL, 0.5 + fwL, lf));
+  }
+
+  // Rounded ground and the seam against loose grains.
+  if (F.r > 0.3) {
+    var m: vec4f = vec4f(isT(n0), isT(n1), isT(n2), isT(n3)) * wq;
+    var sm: f32 = m.x + m.y + m.z + m.w;
+    var blur: vec3f = bsp(tCol, q).rgb; var tc: vec3f = own.rgb;
+    var bd: f32 = 1e9;
+    if (m.x > 0.0) { var e: vec3f = k0.rgb - blur; var dd: f32 = dot(e, e); if (dd < bd) { bd = dd; tc = k0.rgb; } }
+    if (m.y > 0.0) { var e: vec3f = k1.rgb - blur; var dd: f32 = dot(e, e); if (dd < bd) { bd = dd; tc = k1.rgb; } }
+    if (m.z > 0.0) { var e: vec3f = k2.rgb - blur; var dd: f32 = dot(e, e); if (dd < bd) { bd = dd; tc = k2.rgb; } }
+    if (m.w > 0.0) { var e: vec3f = k3.rgb - blur; var dd: f32 = dot(e, e); if (dd < bd) { bd = dd; tc = k3.rgb; } }
+    if (sm < 1e-4) { tc = own.rgb; }
+    tc *= 0.84 + 0.3 * fbm(w * 0.28);
+    tc *= 0.95 + 0.1 * vn(w * 1.6);
+    var gl: f32 = length(gT);
+    var nrm: vec2f = select( vec2f(0.0), -gT / gl , gl > 1e-6 );
+    var band: f32 = 1.0 - smoothstep(0.5, 0.9, M);
+    var lit: f32 = dot(nrm, normalize(vec2f(-0.5, -1.0)));
+    tc *= 1.0 + 0.42 * lit * band;
+    
+    var dm: f32 = (M - 0.5) / fwM;
+    tc *= mix(0.4, 1.0, smoothstep(0.5, 1.4 + zp * 0.08, dm));
+    var dt: f32 = (F.r - 0.5) / fwT;
+    if (nc == 1 && pw > 0.0) { tc *= mix(0.65, 1.0, smoothstep(0.0, 1.5, dt)); }
+    C = mix(C, tc, smoothstep(-0.6, 0.6, (F.r - 0.5) / fwT));
+  }
+
+${GPU_SMOOTH_SLOPE_DRAW}
+  // Round loose grains.
+  var np: i32 = i32(isN(n0, 2) + isN(n1, 2) + isN(n2, 2) + isN(n3, 2));
+  var closed: bool = isM(n0) + isM(n1) + isM(n2) + isM(n3) > 3.5;
+  if (np > 0) {
+    var best: f32 = 1e9; var gc: vec3f = vec3f(0.0); var gd: vec2f = vec2f(0.0); var gr: f32 = 0.6;
+    for (var k: i32 = 0; k < 4; k++) {
+      var o4: vec2i = vec2i(k & 1, k >> 1);
+      var nk: i32 = select(select(select( n3, n2 , k == 2 ), n1 , k == 1 ), n0 , k == 0 );
+      if (nk != 2) { continue; }
+      var ck: vec4f = select(select(select( k3, k2 , k == 2 ), k1 , k == 1 ), k0 , k == 0 );
+      var ce: vec2f = vec2f(c0 + o4) + 0.5;
+      var r: f32 = 0.6 + 0.1 * (h21(ce + (u.view.xy)) - 0.5);
+      var dv: vec2f = (q - ce) / r;
+      var dd: f32 = length(dv);
+      if (dd < best) { best = dd; gc = ck.rgb; gd = dv; gr = r; }
+    }
+    var cover: f32 = select( 1.0 - smoothstep(1.0 - px / gr, 1.0 + px / gr, best), 1.0 , np == 4 );
+    if (closed && F.r < 0.5) { cover = 1.0; }
+    
+    cover *= 1.0 - smoothstep(-0.6, 0.6, (F.r - 0.5) / fwT);
+    if (cover > 0.0) {
+      var r2: f32 = min(1.0, dot(gd, gd));
+      var n3: vec3f = normalize(vec3f(gd * 0.85, sqrt(1.0 - r2) + 0.25));
+      var sh: f32 = 0.98 + 0.18 * max(0.0, dot(n3, normalize(vec3f(-0.45, -0.8, 0.6))));
+      sh *= mix(1.0, 0.94, smoothstep(0.7 - px * 0.2, 1.15 + px * 0.2, best));
+      var pc: vec3f = (select( gc, (k0.rgb * wq.x + k1.rgb * wq.y + k2.rgb * wq.z + k3.rgb * wq.w) ,np == 4 )) * sh;
+      if (!closed) { pc *= mix(1.0, 0.72, smoothstep(0.7, 1.0, best)); }
+      C = mix(C, pc, cover);
+    }
+  }
+
+  // Built pieces: belts, screws, wall tiles and filters.
+  if (nc == 4) { C = own.rgb; }
+  else if (nc == 7 || nc == 8) {
+    
+    var dir: f32 = select( -1.0, 1.0 , nc == 8 );
+    var lb: vec2f = fract(w * 0.25) * 4.0;
+    
+    var sd: f32 = beltLineSD(c, lb);
+    var inside: f32 = 1.0 - smoothstep(-px, px, sd);
+    if (inside > 0.0) {
+      var edge: f32 = -sd;
+      var bc: vec3f;
+      if (edge < 0.8) {
+        bc = bandColor(lb, bandRidge(w, lb, dir));
+        bc *= mix(0.55, 1.0, smoothstep(0.0, max(px * 1.4, 0.12), edge));
+      } else {
+        var rc: vec2f = vec2f(select( 3.0, 1.0 ,lb.x < 2.0 ), 2.0); var dv: vec2f = lb - rc;
+        var rd: f32 = length(dv);
+        bc = vec3f(0.15, 0.16, 0.18);
+        if (rd < 0.95) {
+          var ang: f32 = atan2(dv.y, dv.x) - dir * (su.x) * 30.0 / 0.95;
+          var spoke: f32 = smoothstep(0.55, 0.85, cos(ang * 3.0));
+          bc = mix(vec3f(0.5, 0.36, 0.25), vec3f(0.72, 0.53, 0.36), 0.5 + 0.5 * dot(normalize(dv + 1e-4), vec2f(-0.6, -0.8)) * smoothstep(0.2, 0.95, rd));
+          bc *= mix(1.0, 0.72, spoke * smoothstep(0.25, 0.4, rd));
+          bc *= mix(1.0, 0.55, smoothstep(0.82, 0.95, rd));
+        }
+      }
+      C = mix(C, bc, inside);
+    }
+  }
+  else if ((nc >= 9 && nc <= 12) || (nc >= 19 && nc <= 22)) {
+    
+    var vert: bool = nc == 9 || nc == 11;
+    var sp2: f32 = select(select(select( -1.0, 1.0 , nc == 11 ), 1.0 , nc == 10 ), -1.0 , nc == 9 );
+    var a: f32 = (select( fract(w.y * 0.25), fract(w.x * 0.25) ,vert )) * 4.0; var along: f32 = select( w.x, w.y , vert );
+    var screwSD: f32 = -1.0; var capRing: f32 = 0.0;
+    var screwTag: vec4f = cellAt(tTree, c);
+    var screwJoins: i32 = i32(screwTag.a * 255.0 + 0.5);
+    
+    var shaped: bool = nc >= 19 || (screwJoins > 0 && screwJoins < 64);
+    if (shaped) {
+      var d: vec2f = vec2f(1.0,select( 1.0, -1.0 , nc == 19 || nc == 21 ));
+      var tag: vec4f = screwTag;
+      var local: vec2f = floor(tag.gb * 255.0 + 0.5) - 8.0 + fract(w) - 2.0;
+      var joins: i32 = screwJoins;
+      screwSD = 100.0;
+      for (var end: i32 = 0; end < 2; end++) {
+        var join: i32 = (joins >> u32(end * 3)) & 7;
+        var side: f32 = select( 1.0, -1.0 , end == 0 );
+        var diagonal: bool = select( join != 0, join == 0 || join >= 5 , nc >= 19 );
+        var grade: f32 = select(select( 1.0, -1.0 , join == 1 ), d.y , nc >= 19 );
+        var ray: vec2f = select(select(select(select(select( vec2f(-1, 0), vec2f(0, 1) , join == 3 ), vec2f(1, 0) , join == 2 ), vec2f(0, -1) ,
+          join == 1 ), (select( vec2f(side, 0), vec2f(0, side) ,vert )) ,
+          nc < 19 ), (select( vec2f(side, side * grade), vec2f(side * grade, side) ,vert )) * 0.70710678 , diagonal );
+        var t: f32 = dot(local, ray); var cross2: f32 = dot(local, vec2f(-ray.y, ray.x));
+        var radius: f32 = length(vec2f(cross2, min(0.0, t))); var cap: f32 = t - (select( 2.0, 2.82842712 ,diagonal ));
+        
+        var sd: f32 = max(radius - 2.0,select( -10.0, cap , (joins & (64 << u32(end))) != 0 ));
+        if (sd < screwSD) {
+          screwSD = sd;
+          a = 2.0 + (select( radius, -radius ,side * cross2 < 0.0 ));
+          along = (select(select( w.x - local.x, w.y - local.y , vert ), dot(w - local, d) * 0.70710678 ,nc >= 19 )) + side * t;
+          capRing = select( 0.0, smoothstep(-0.55, -0.15, cap) ,(joins & (64 << u32(end))) != 0 );
+        }
+      }
+      if (nc >= 19) { sp2 = select( -1.0, 1.0 ,nc == 19 || nc == 20 ); }
+    }
+    var wall: f32 = 0.62;
+    var tc: vec3f;
+    if (a < wall || a > 4.0 - wall || capRing > 0.5) {
+      var u2: f32 = select( (4.0 - a) / wall, a / wall , a < wall );   
+      var cyl: f32 = sin(clamp(u2, 0.0, 1.0) * 1.5708);
+      tc = vec3f(0.62, 0.41, 0.27) * (0.6 + 0.55 * cyl);
+      if (a < wall) { tc *= 1.08; }
+      var ring: f32 = 1.0 - smoothstep(0.0, 0.3, min(fract(along * 0.25), 1.0 - fract(along * 0.25)) * 4.0);
+      if (nc < 19) { tc *= 1.0 - 0.3 * ring; }
+      tc *= mix(0.5, 1.0, smoothstep(0.0, max(px * 1.3, 0.1), min(a, 4.0 - a)));
+    } else {
+      var full: bool = own.r + own.g + own.b > 0.0;
+      var fl: f32 = fract((along + a * 0.9 - sp2 * (su.x) * 14.0) * 0.5);
+      var flight: f32 = smoothstep(0.82, 0.92, fl) * (1.0 - smoothstep(0.92, 1.0, fl));
+      var bore: vec3f = vec3f(0.11, 0.1, 0.11) * (0.8 + 0.4 * (1.0 - abs(a - 2.0) / 1.4));
+      tc = select( bore, own.rgb * (0.9 + 0.1 * vn(w * 1.3)) ,full );
+      tc = mix(tc, vec3f(0.6, 0.48, 0.38), flight * (select( 0.8, 0.35 ,full )));
+      tc *= 1.0 - 0.25 * smoothstep(0.9, 1.3, abs(a - 2.0));
+    }
+    if (shaped) {
+      var full: bool = own.r + own.g + own.b > 0.0;
+      if (full) { C = mix(C, own.rgb, 1.0 - smoothstep(0.36, 0.5, length(fract(w) - 0.5))); }
+      C = mix(C, tc, 1.0 - smoothstep(2.0 - px, 2.0 + px, screwSD + 2.0));
+    } else { C = tc; }
+  }
+  else if (nc == 6 || (nc >= 15 && nc <= 18)) {
+    var lb: vec2f = fract(w * 0.25) * 4.0; var rb: vec2f = 4.0 - lb;
+    var e: f32 = min(min(lb.x, lb.y), min(rb.x, rb.y));
+    var tc: vec3f = own.rgb * (0.92 + 0.14 * fbm(w * 0.55));
+    var hl: f32 = max(1.0 - smoothstep(0.0, 0.8, lb.x), 1.0 - smoothstep(0.0, 0.8, lb.y));
+    var shd: f32 = max(1.0 - smoothstep(0.0, 0.8, rb.x), 1.0 - smoothstep(0.0, 0.8, rb.y));
+    tc *= 1.0 + 0.2 * hl - 0.26 * shd;
+    tc *= mix(0.5, 1.0, smoothstep(px * 0.5, px * 0.5 + max(px * 1.2, 0.14), e));
+    if (nc >= 15) {
+      var sl: vec2f = lb;
+      if (nc == 16 || nc == 18) { sl.x = 4.0 - sl.x; }
+      if (nc >= 17) { sl.y = 4.0 - sl.y; }
+      var d: f32 = (sl.y - sl.x) * 0.70710678;
+      tc *= 1.0 + 0.2 * (1.0 - smoothstep(0.0, 0.6, d));
+      C = mix(C, tc, smoothstep(-px * 0.5, px * 0.5, d));
+    } else { C = tc; }
+  }
+
+  else if (nc == 13 || nc == 14) {
+
+    var dir: f32 = select( -1.0, 1.0 , nc == 14 );
+    var lb: vec2f = fract(w * 0.25) * 4.0;
+    var tag: vec4f = cellAt(tTree, c);
+    var carry: bool = tag.a > 0.16; var tinted: bool = tag.r + tag.g + tag.b > 0.0;
+    var sd: f32 = beltLineSD(c, lb);
+    var inside: f32 = 1.0 - smoothstep(-px, px, sd);
+    if (inside > 0.0) {
+      var edge: f32 = -sd;
+      var steel: vec3f = vec3f(0.42, 0.44, 0.5);
+      
+      var behind: vec3f = select( C * 0.8, own.rgb , own.r + own.g + own.b > 0.0 );
+      var cg: vec2f = abs(fract(w) - 0.5);
+      var bw2: f32 = max(px * 0.9, 0.09);
+      var hole: f32 = 1.0 - smoothstep(0.5 - bw2 - px, 0.5 - bw2, max(cg.x, cg.y));
+      var fc: vec3f;
+      if (edge < 0.8) {
+        var rg: f32 = bandRidge(w, lb, dir);
+        var band: vec3f = bandColor(lb, rg);
+        if (tinted && (carry == (lb.y < 2.0))) { band = mix(band, tag.rgb * (select( 0.8, 1.0 ,lb.y < 2.0 )), 0.75 * rg); }
+        fc = mix(band, behind * 0.85, hole * (1.0 - rg) * 0.75);
+        fc *= mix(0.55, 1.0, smoothstep(0.0, max(px * 1.4, 0.12), edge));
+      } else {
+        fc = mix(steel * 0.78, behind, hole);
+        fc *= mix(0.7, 1.0, smoothstep(0.0, max(px * 1.4, 0.12), edge - 0.8));
+      }
+      C = mix(C, fc, inside);
+    }
+  }
+
+  // Light, then the legacy nine-tap glow filter.
+  if (uLightMap.w > 0.5) { C *= sampleLevel(tLight, (q - uLightMap.xy) / uLightMap.z / (vec2f(textureDimensions(tLight))), 0.0).rgb; }
+  if ((su.z) > 0.0) {
+    var gu: vec2f = q / (2.0 * (vec2f(textureDimensions(tGlow)))); var d: vec2f = 1.0 / (vec2f(textureDimensions(tGlow)));
+    var g: vec3f = sampleLevel(tGlow, gu, 0.0).rgb * 0.25;
+    g += (sampleLevel(tGlow, gu + vec2f(d.x, 0), 0.0).rgb + sampleLevel(tGlow, gu - vec2f(d.x, 0), 0.0).rgb + sampleLevel(tGlow, gu + vec2f(0, d.y), 0.0).rgb + sampleLevel(tGlow, gu - vec2f(0, d.y), 0.0).rgb) * 0.125;
+    g += (sampleLevel(tGlow, gu + d, 0.0).rgb + sampleLevel(tGlow, gu - d, 0.0).rgb + sampleLevel(tGlow, gu + vec2f(d.x, -d.y), 0.0).rgb + sampleLevel(tGlow, gu + vec2f(-d.x, d.y), 0.0).rgb) * 0.0625;
+    C += g * (su.z);
+  }
+  return vec4f(C, 1.0);
+}`;

src/render/gpu/smooth-slope.js

@@ -0,0 +1,107 @@
+// Smooth slope geometry and cleats, ported from slope-belt-shader.js.
+const GPU_SMOOTH_SLOPE_HELPERS = `fn slopeStroke(p: vec2f, a: vec2f, b: vec2f, origin: f32) -> vec3f {
+  var ray: vec2f = b - a; var len: f32 = length(ray); ray /= len;
+  var along: f32 = dot(p - a, ray); var across: f32 = dot(p - a, vec2f(-ray.y, ray.x));
+  var q: vec2f = abs(vec2f(along - len * 0.5, across)) - vec2f(len * 0.5 - 0.75, 1.1);
+  var sd: f32 = length(max(q, vec2f(0.0))) + min(max(q.x, q.y), 0.0) - 0.75;
+  return vec3f(sd, origin + along, across + 2.0);
+}
+fn slopeCorner(p: vec2f, a: vec2f, origin: f32) -> vec3f {
+  var d: vec2f = p - a; var radius: f32 = length(d);
+  return vec3f(radius - 1.85, origin + d.x, 2.0 + d.y);
+}
+fn slopeShape(inputP: vec2f, worldOrigin: vec2f, grade: f32, joins: i32) -> vec3f { var p=inputP;
+  if ((joins & 64) != 0) { p.y -= 0.45; }
+  var sy: f32 = select( 0.0, 4.0 , grade < 0.0 );
+  var a: vec2f = vec2f(-grade * 1.30814755, sy + 1.30814755); var b: vec2f = vec2f(4.0 - grade * 1.30814755, 4.0 - sy + 1.30814755);
+  if ((joins & 1) != 0) { a.y = sy + 2.0; a.x = (a.y - sy - 2.61629509) / grade; }
+  else if ((joins & 4) != 0) { a -= vec2f(4.0, grade * 4.0); }
+  if ((joins & 2) != 0) { b.y = 6.0 - sy; b.x = (b.y - sy - 2.61629509) / grade; }
+  else if ((joins & 8) != 0) { b += vec2f(4.0, grade * 4.0); }
+  var axis: vec2f = vec2f(1.0, grade) * 0.70710678;
+  var best: vec3f = slopeStroke(p, a, b, dot(worldOrigin + a, axis));
+  if ((joins & 1) != 0) {
+    var h: vec3f = slopeStroke(p, vec2f(-2.0, a.y), a, worldOrigin.x - 2.0);
+    if (h.x < best.x) { best = h; }
+    h = slopeCorner(p, a, worldOrigin.x + a.x);
+    if (h.x < best.x) { best = h; }
+  }
+  if ((joins & 2) != 0) {
+    var h: vec3f = slopeStroke(p, b, vec2f(6, b.y), worldOrigin.x + b.x);
+    if (h.x < best.x) { best = h; }
+    h = slopeCorner(p, b, worldOrigin.x + b.x);
+    if (h.x < best.x) { best = h; }
+  }
+  return best;
+}
+
+`;
+const GPU_SMOOTH_SLOPE_DRAW = `  if (nc == 31 || nc == 32) {
+    var tag: vec4f = cellAt(tTree, c);
+    var grade: f32 = select( 1.0, -1.0 , nc == 31 );
+    var local: vec2f = floor(tag.gb * 255.0 + 0.5) - 8.0 + fract(w);
+    var sy: f32 = select( 0.0, 4.0 , grade < 0.0 );
+    var foot: f32 = (local.x + grade * (local.y - sy)) * 0.5;
+    var depth: f32 = (grade * local.x - local.y + sy) * 0.70710678;
+    var wallFlags: i32 = i32(tag.a * 255.0 + 0.5);
+    var ends: f32 = max(select( -100.0, -foot ,(wallFlags & 8) == 0 ),select( -100.0, foot - 4.0 , (wallFlags & 16) == 0 ));
+    var sd: f32 = max(ends * 1.41421356, max(-depth, depth - ${CLEAT_HEIGHT.toFixed(1)}));
+    var phase: f32 = f32(wallFlags & 7);
+    var wave: f32 = 0.5 + 0.5 * sin((foot + w.x - local.x - phase) * 3.14159265);
+    var rubber: vec3f = mix(vec3f(0.20,0.22,0.25), vec3f(0.25,0.27,0.31), wave);
+    C = mix(C, rubber, 1.0 - smoothstep(-px, px, sd));
+
+    var dirS: f32 = select( -1.0, 1.0 , (wallFlags & 32) != 0 );
+    if (grade * dirS < 0.0) {
+      var rc: f32 = (w.x - local.x) + foot - (select( 0.0, 1.0 ,grade < 0.0 )) - dirS * phase;
+      var off: f32 = (modulo(rc + ${(CLEAT_SPACING / 2).toFixed(1)}, ${CLEAT_SPACING.toFixed(1)}) - ${(CLEAT_SPACING / 2).toFixed(1)}) * 1.41421356;
+      var ribTop: f32 = ${((2 * CLEAT_REACH + 1) * Math.SQRT1_2).toFixed(3)};
+      var rd: f32 = length(vec2f(off, depth - clamp(depth, 0.0, ribTop))) - 0.5;
+      var rib: vec3f = mix(vec3f(0.50,0.52,0.56), vec3f(0.70,0.72,0.76), smoothstep(ribTop - 1.2, ribTop, depth));
+      rib *= mix(0.78, 1.0, smoothstep(-0.5, 0.2, off * -dirS * grade));
+      C = mix(C, rib, 1.0 - smoothstep(-px, px, rd));
+    }
+  }
+
+  if ((nc == 2 || (nc >= 23 && nc <= 30)) &&
+      (i32(cellAt(tTree, c).a * 255.0 + 0.5) & 128) != 0 &&
+      cellAt(tTree, c).r * 255.0 >= 22.5 && cellAt(tTree, c).r * 255.0 <= 30.5) {
+    var slopeClass: i32 = i32(cellAt(tTree, c).r * 255.0 + 0.5);
+    var shape: i32 = (slopeClass - 23) / 2;
+    var mesh: bool = shape >= 2;
+    var grade: f32 = select( 1.0, -1.0 , (shape == 0 || shape == 2) );
+    var dir: f32 = select( -1.0, 1.0 , ((slopeClass - 23) % 2) == 1 );
+    var tag: vec4f = cellAt(tTree, c);
+    var joins: i32 = i32(tag.a * 255.0 + 0.5);
+    var local: vec2f = floor(tag.gb * 255.0 + 0.5) - 8.0 + fract(w);
+    var shapePos: vec3f = slopeShape(local, w - local, grade, joins);
+    var sd: f32 = shapePos.x; var axis: f32 = shapePos.y; var depth: f32 = shapePos.z;
+    var time: f32 = select( 0.0, (su.x) , !mesh && (joins & 16) != 0 );
+    var lb: vec2f = vec2f(axis, depth); var alongWorld: vec2f = vec2f(axis, depth);
+    var rg: f32 = step(0.72, fract((axis - dir * (select( -1.0, 1.0 ,depth < 2.0 )) * time * 30.0) * 0.5));
+    var belt: vec3f;
+    if (-sd < select(0.8,1.1,mesh)) {
+      belt = select( bandColor(lb, rg), mix(vec3f(0.22,0.25,0.29), vec3f(0.48,0.53,0.59),select( 0.35, 0.85 , depth < 2.0 )) ,mesh );
+      belt *= mix(0.55, 1.0, smoothstep(0.0, max(px * 1.4, 0.12), -sd));
+    } else if (mesh) {
+      var cg: vec2f = abs(fract(alongWorld) - 0.5);
+      var bw2: f32 = max(px * 0.9, 0.09);
+      var hole: f32 = 1.0 - smoothstep(0.5 - bw2 - px, 0.5 - bw2, max(cg.x, cg.y));
+      var steel: vec3f = select( vec3f(0.42, 0.44, 0.5), mix(vec3f(0.42, 0.44, 0.5), own.rgb, 0.4) , dot(own.rgb, vec3f(1)) > 0.01 );
+      belt = mix(steel * 0.78, C * 0.8, hole);
+    } else {
+      var dv: vec2f = vec2f(modulo(axis, 2.0) - 1.0, depth - 2.0);
+      var rd: f32 = length(dv);
+      belt = vec3f(0.15, 0.16, 0.18);
+      if (rd < 0.95) {
+        var ang: f32 = atan2(dv.y, dv.x) - dir * time * 30.0 / 0.95;
+        var spoke: f32 = smoothstep(0.55, 0.85, cos(ang * 3.0));
+        belt = mix(vec3f(0.5, 0.36, 0.25), vec3f(0.72, 0.53, 0.36), 0.5 + 0.5 * dot(normalize(dv + 1e-4), vec2f(-0.6, -0.8)) * smoothstep(0.2, 0.95, rd));
+        belt *= mix(1.0, 0.72, spoke * smoothstep(0.25, 0.4, rd));
+        belt *= mix(1.0, 0.55, smoothstep(0.82, 0.95, rd));
+      }
+    }
+    C = mix(C, belt, 1.0 - smoothstep(-px, px, sd));
+  }
+
+`;

src/render/gpu/smooth.js

@@ -0,0 +1,47 @@
+// Field mip levels preserve the legacy trilinear AO sample at level 2.2.
+const GPU_FIELD_MIP = `
+@group(0) @binding(0) var source: texture_2d<f32>;
+@group(0) @binding(1) var linear: sampler;
+@vertex fn vs(@builtin(vertex_index) i: u32) -> @builtin(position) vec4f {
+  let pos=array<vec2f,3>(vec2f(-1,-1),vec2f(3,-1),vec2f(-1,3)); return vec4f(pos[i],0,1);
+}
+@fragment fn fs(@builtin(position) p: vec4f) -> @location(0) vec4f {
+  let size=max(vec2u(1),textureDimensions(source)/2u);
+  return textureSampleLevel(source,linear,p.xy/vec2f(size),0);
+}`;
+function gpuSmoothInit() {
+  const d=GPU.device;
+  const pipeline=(code,vertex,format)=>d.createRenderPipeline({layout:'auto',vertex:{module:d.createShaderModule({code:vertex}),entryPoint:'vs'},fragment:{module:d.createShaderModule({code}),entryPoint:'fs',targets:[{format}]},primitive:{topology:'triangle-list'}});
+  GPU.smooth=pipeline(GPU_SMOOTH,GPU_DRAW,GPU.format);
+  GPU.fieldMip=pipeline(GPU_FIELD_MIP,GPU_FIELD_MIP,'rgba8unorm');
+  GPU.smoothUniform=d.createBuffer({size:16,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST});
+}
+function gpuSmoothTextures() {
+  const w=R.bw,h=R.bh;
+  GPU.smoothOver=Array.from({length:3},()=>gpuTexture(w,h));
+  GPU.smoothMask=gpuTexture(w,h,'r8uint');
+  GPU.backTex=gpuTexture(w,h,'rgba8unorm',true);
+  GPU.treeTex=gpuTexture(w,h,'rgba8unorm',true);
+  GPU.fieldTex=GPU.device.createTexture({size:[w,h],format:'rgba8unorm',mipLevelCount:Math.min(4,1+Math.floor(Math.log2(Math.max(w,h)))),usage:GPUTextureUsage.TEXTURE_BINDING|GPUTextureUsage.STORAGE_BINDING|GPUTextureUsage.RENDER_ATTACHMENT});
+  GPU.viewTextures.push(...GPU.smoothOver,GPU.smoothMask,GPU.backTex,GPU.treeTex,GPU.fieldTex);
+}
+function gpuSmoothBindings(group,view,buffer) {
+  GPU.bindColourSmooth=group(GPU.colourSmooth,[buffer(GPU.uniform),buffer(GPU.lut),buffer(GPU.land),...GPU.worldTextures.slice(0,9).map(view),view(GPU.override),view(GPU.overrideGlow),view(GPU.colourTex),view(GPU.glowTex),view(GPU.backTex),view(GPU.treeTex),GPU.fieldTex.createView({baseMipLevel:0,mipLevelCount:1}),...GPU.smoothOver.map(view),view(GPU.smoothMask)]);
+  GPU.bindSmooth=group(GPU.smooth,[buffer(GPU.uniform),view(GPU.colourTex),view(GPU.backTex),view(GPU.treeTex),view(GPU.fieldTex),view(GPU.lightTex),view(GPU.glowTex),GPU.sampler,buffer(GPU.smoothUniform)]);
+  GPU.fieldMips=[];
+  for(let i=1;i<GPU.fieldTex.mipLevelCount;i++) GPU.fieldMips.push({
+    bind:group(GPU.fieldMip,[GPU.fieldTex.createView({baseMipLevel:i-1,mipLevelCount:1}),GPU.sampler]),
+    view:GPU.fieldTex.createView({baseMipLevel:i,mipLevelCount:1})
+  });
+}
+function gpuSmoothUpload() {
+  for(const [i,a] of [R.backPx,R.treePx,R.fieldPx].entries()) gpuWrite(GPU.smoothOver[i],a,R.bw,0,0,R.bw,R.bh);
+  gpuWrite(GPU.smoothMask,GPU.overrideMask,R.bw,0,0,R.bw,R.bh);
+  GPU.device.queue.writeBuffer(GPU.smoothUniform,0,new Float32Array([(performance.now()-SM.t0)/1000,smoothSurface(),SET.glow==='off'?0:SET.glow==='high'?0.9:0.55,0]));
+}
+function gpuSmoothMips(encoder) {
+  for(const m of GPU.fieldMips) {
+    const pass=encoder.beginRenderPass({colorAttachments:[{view:m.view,loadOp:'clear',storeOp:'store'}]});
+    pass.setPipeline(GPU.fieldMip);pass.setBindGroup(0,m.bind);pass.draw(3);pass.end();
+  }
+}

src/render/gpu/textures.js

@@ -35,9 +35,9 @@ function gpuWorldTextures() {
   GPU.bindColour = null;
 }
 function gpuViewTextures() {
-  const key = `${R.bw}/${R.bh}/${R.lw}/${R.lh}`;
+  const key = `${R.bw}/${R.bh}/${R.lw}/${R.lh}/${+R.smoothOn}`;
   if (GPU.viewKey === key) return;
-  GPU.viewKey = key;
+  GPU.viewKey = key; GPU.lightFrame = undefined;
   if (GPU.viewTextures) for (const t of GPU.viewTextures) t.destroy();
   const bw = R.bw, bh = R.bh, gw = R.glowC.width, gh = R.glowC.height;
   GPU.override = gpuTexture(bw, bh); GPU.overrideGlow = gpuTexture(gw, gh);
@@ -47,6 +47,7 @@ function gpuViewTextures() {
   GPU.lightTex = gpuTexture(Math.max(1, R.lw), Math.max(1, R.lh), 'rgba8unorm', true);
   GPU.lightWork = [0,1].map(() => gpuTexture(Math.max(1,R.lw),Math.max(1,R.lh),'rgba16float',true));
   GPU.viewTextures = [GPU.override, GPU.overrideGlow, GPU.colourTex, GPU.glowTex, GPU.blurX, GPU.blurY, GPU.lightTex, ...GPU.lightWork];
+  if (R.smoothOn) gpuSmoothTextures();
   GPU.bindColour = null;
 }
 function gpuBindings() {
@@ -54,6 +55,7 @@ function gpuBindings() {
   const d = GPU.device, buffer = b => ({ buffer: b }), view = t => t.createView();
   const group = (pipeline, resources) => d.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: resources.map((resource, binding) => ({ binding, resource })) });
   GPU.bindColour = group(GPU.colour, [buffer(GPU.uniform), buffer(GPU.lut), buffer(GPU.land), ...GPU.worldTextures.slice(0,9).map(view), view(GPU.override), view(GPU.overrideGlow), view(GPU.colourTex), view(GPU.glowTex)]);
+  if (R.smoothOn) gpuSmoothBindings(group, view, buffer);
   gpuLightBindings(group, view, buffer);
   GPU.bindDraw = group(GPU.draw, [buffer(GPU.uniform), view(GPU.colourTex), view(GPU.lightTex), view(GPU.blurY), GPU.sampler]);
   if (GPU.blurUniforms) for (const b of GPU.blurUniforms) b.destroy();
@@ -89,6 +91,7 @@ function gpuUpload(lighting, glowOn) {
   gpuWrite(GPU.overrideGlow, R.glowPx, R.glowC.width, 0, 0, R.glowC.width, R.glowC.height);
   for(let i=0;i<W;i++) { GPU.columnData[i*4]=surfY[i]; GPU.columnData[i*4+1]=roofY[i]; }
   q.writeBuffer(GPU.columns,0,GPU.columnData);
+  if (R.smoothOn) gpuSmoothUpload();
   const zp = R.zp, ox = Math.round(-(R.camX - R.vx0) * zp), oy = Math.round(-(R.camY - R.vy0) * zp);
   GPU.params.set([R.vx0, R.vy0, R.bw, R.bh, W, H, FW, FH,
     +SET.fog, R.rain, SET.grain, R.frame, +R.lens, Math.max(0, Object.keys(THERM).indexOf(SET.thermalPalette)), +SET.incandescence, +glowOn,

src/render/objects.js

@@ -16,7 +16,7 @@ function worldObjects() {
       if (dy === 0) { r += end ? 10 : 34; g += end ? 8 : 28; b += end ? 6 : 24; } else { r -= end ? 34 : 18; g -= end ? 30 : 16; b -= end ? 26 : 14; }
       if (e.T > 450 && incand) { const gi = glowIndex(e.T), w = GLOW_W[gi]; r = r * (1 - w) + GLOW_R[gi] * w; g = g * (1 - w) + GLOW_G[gi] * w; b = b * (1 - w) + GLOW_B[gi] * w; }
       if (r < 0) r = 0; if (g < 0) g = 0; if (b < 0) b = 0;
-      if (R.gpuOn && !px[sy * bw + sx]) GPU.overrideCount++;
+      if (R.gpuOn && !gpuOverrideHas(sy * bw + sx)) { if (sm) cellPixel(sx + vx0, sy + vy0); else GPU.overrideCount++; }
       px[sy * bw + sx] = (sm ? 0x04000000 : 0xff000000) | ((b > 255 ? 255 : b | 0) << 16) | ((g > 255 ? 255 : g | 0) << 8) | (r > 255 ? 255 : r | 0);
     }
   }
@@ -26,7 +26,7 @@ function worldObjects() {
     const sx = Math.floor(f.x) - vx0, sy = Math.floor(f.y) - vy0;
     if (sx < 0 || sy < 0 || sx >= bw || sy >= bh) continue;
     const m = f.m, v = ((sx * 7 + sy * 13) & 15) - 8;
-    if (R.gpuOn && !px[sy * bw + sx]) GPU.overrideCount++;
+    if (R.gpuOn && !gpuOverrideHas(sy * bw + sx)) { if (sm) cellPixel(sx + vx0, sy + vy0); else GPU.overrideCount++; }
     px[sy * bw + sx] = (sm ? 0x02000000 : 0xff000000) | (Math.max(0, Math.min(255, COLB[m] + v)) << 16) | (Math.max(0, Math.min(255, COLG[m] + v)) << 8) | Math.max(0, Math.min(255, COLR[m] + v));
   }
   // blueprint ghosts: blocks waiting to be built
@@ -68,7 +68,7 @@ function worldObjects() {
         const sx = gbx * BS + dx - vx0, sy = gby * BS + dy - vy0;
         if (sx < 0 || sy < 0 || sx >= bw || sy >= bh) continue;
         const k = sy * bw + sx;
-        if (R.gpuOn && !px[k]) cellPixel(sx + vx0, sy + vy0);
+        if (R.gpuOn && !gpuOverrideHas(k)) cellPixel(sx + vx0, sy + vy0);
         const c = px[k];
         const edge = slope || diagonal ? sd > -0.8 : (dx === 0 && (caps & 1)) || dy === 0 || (dx === BS - 1 && (caps & 2)) || dy === BS - 1;
         const a = edge ? (((dx + dy + pulse) & 3) === 0 ? 0.85 : 0.6) : 0.28;

src/render/smooth.js

@@ -55,12 +55,17 @@ function smUpload(name, w, h, data) {
   if (T.mip) gl.generateMipmap(gl.TEXTURE_2D);
 }
 const smBytes = a => new Uint8Array(a.buffer, a.byteOffset, a.byteLength);
+// Both renderers share the same cached parallax horizon.
+function smoothSurface() {
+  if (SM.surf < 0 || SM.surfSeed !== worldSeed || SM.surfW !== W) { SM.surf = meanSurface(); SM.surfW = W; SM.surfSeed = worldSeed; }
+  return SM.surf;
+}
 function smoothDraw(lighting, glowOn) {
   const t0 = performance.now();
   // the GPU pass can run below full resolution on a slow graphics card; it is scaled up smoothly
   const gl = SM.gl, cv = SM.canvas, sc = SM.scale, cw = Math.max(1, Math.round(R.canvas.width * sc)), ch = Math.max(1, Math.round(R.canvas.height * sc)), u = SM.u;
   if (cv.width !== cw || cv.height !== ch) { cv.width = cw; cv.height = ch; }
-  if (SM.surf < 0 || SM.surfSeed !== worldSeed || SM.surfW !== W) { SM.surf = meanSurface(); SM.surfW = W; SM.surfSeed = worldSeed; }
+  const surf = smoothSurface();
   gl.viewport(0, 0, cw, ch);
   gl.useProgram(SM.prog);
   const bw = R.bw, bh = R.bh;
@@ -82,7 +87,7 @@ function smoothDraw(lighting, glowOn) {
   gl.uniform2f(u.uGlowSize, R.glowC.width || 1, R.glowC.height || 1);
   gl.uniform1f(u.uGlowK, glowOn ? (SET.glow === 'high' ? 0.9 : 0.55) : 0);
   gl.uniform1f(u.uTime, (performance.now() - SM.t0) / 1000);
-  gl.uniform1f(u.uSurf, SM.surf);
+  gl.uniform1f(u.uSurf, surf);
   gl.uniform1f(u.uRain, R.rain || 0);
   gl.drawArrays(gl.TRIANGLES, 0, 3);
   const ctx = R.ctx;

src/render/world.js

@@ -6,7 +6,7 @@ function drawWorld() {
   const perfStart = (PERF_MARKS || SET.showPerf) ? performance.now() : 0;
   if (PERF_MARKS || SET.showPerf) { perfMarks.light = 0; perfMarks.override = 0; delete perfMarks.gpu; }
   const wasGpu = R.gpuOn;
-  R.gpuOn = SET.renderer === 'webgpu' && SET.look === 'pixel' && GPU.ok;
+  R.gpuOn = SET.renderer === 'webgpu' && GPU.ok;
   if (R.gpuOn && !wasGpu) GPU.lightFrame = undefined;
   GPU.canvas.hidden = !R.gpuOn;
   if (R.gpuOn) R.ctx.clearRect(0, 0, R.canvas.width, R.canvas.height);
@@ -15,9 +15,8 @@ function drawWorld() {
   const vx0 = R.vx0, vy0 = R.vy0;
   // how hard it is raining around the camera, eased so the sky greys over gently
   R.rain = (R.rain || 0) * 0.98 + rainAt(Math.max(0, Math.min(W - 1, Math.floor(R.camX + bw / 2)))) * 0.02;
-  // Smooth look: the colours below are still worked out per grain, but each grain also gets a
+  // Smooth inputs carry a class in alpha; only exceptional cells use the CPU in WebGPU mode.
-  // class (alpha byte) and the GPU draws rounded shapes from them (see smooth-shaders.js)
+  const sm = R.smoothOn = SET.look === 'smooth' && !R.lens && (R.gpuOn || smoothAvailable());
-  const sm = R.smoothOn = SET.look === 'smooth' && !R.lens && typeof smoothAvailable === 'function' && smoothAvailable();
   const mmask = R.moldMask;
   if (sm) {
     mmask.fill(0);

src/ui/loop.js

@@ -67,7 +67,7 @@ function frame(now) {
     weather.seen = true;
     showCard({ kind: 'Weather', title: 'Rain', swatch: WATER, body: ['Steam that rises into the open sky gathers into clouds and comes back down as rain nearby. On freezing heights it falls as snow.', 'Water is never used up, only moved around. If some leaves the world (dropped into the dock, say), a shower from beyond the claim makes it up.'] });
   }
-  if (SET.look === 'smooth' && SM.tried && !SM.ok) {
+  if (SET.look === 'smooth' && !R.gpuOn && SM.tried && !SM.ok) {
     SET.look = 'pixel'; saveSettings();
     toast(SM.why === 'slow' ? 'The smooth look was too slow on this computer, so the pixel look is on. You can try it again in Settings.' : SM.why === 'software' ? 'This browser is drawing without the graphics card, so the smooth look would crawl. The pixel look is on.' : 'This browser cannot draw the smooth look, so the pixel look is on.');
   }

src/ui/settings.js

@@ -16,7 +16,7 @@ const SET_TABS = {
     { k: 'shake', label: 'Screen shake', type: 'toggle' },
   ],
   Graphics: [
-    { k: 'renderer', label: 'Renderer', d: 'WebGPU colours the pixel look on the graphics card. Smooth uses the original renderer.', type: 'select', opts: [['webgpu', 'WebGPU'], ['legacy', 'Legacy']] },
+    { k: 'renderer', label: 'Renderer', d: 'WebGPU draws both looks on the graphics card. Legacy uses the original renderer.', type: 'select', opts: [['webgpu', 'WebGPU'], ['legacy', 'Legacy']] },
     { k: 'look', label: 'Look', d: 'Smooth draws rounded, outlined ground, round grains and glossy liquids on the graphics card. Pixel is the original look, one square per grain.', type: 'select', opts: [['smooth', 'Smooth'], ['pixel', 'Pixel']] },
     { k: 'lighting', label: 'Lighting', d: 'Full is colored light at grain resolution. Fast uses half resolution.', type: 'select', opts: [['off', 'Off'], ['fast', 'Fast'], ['full', 'Full']] },
     { k: 'lightEvery', label: 'Light updates', type: 'select', opts: [[1, 'Every frame'], [2, 'Every 2nd frame'], [3, 'Every 3rd frame']] },

test/gpu-readback.js

@@ -0,0 +1,30 @@
+// Read the four actual smooth textures without adding COPY_SRC usage to game resources.
+async function testSmoothInputs() {
+  const d=GPU.device, code=`
+        @group(0) @binding(0) var col: texture_2d<f32>;
+        @group(0) @binding(1) var back: texture_2d<f32>;
+        @group(0) @binding(2) var tree: texture_2d<f32>;
+        @group(0) @binding(3) var field: texture_2d<f32>;
+        @group(0) @binding(4) var<storage,read_write> out: array<u32>;
+        @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) id: vec3u) {
+          let size=textureDimensions(col); let k=id.x; if(k>=size.x*size.y){return;}
+          let p=vec2i(i32(k%size.x),i32(k/size.x));
+          out[k*4]=pack4x8unorm(textureLoad(col,p,0));
+          out[k*4+1]=pack4x8unorm(textureLoad(back,p,0));
+          out[k*4+2]=pack4x8unorm(textureLoad(tree,p,0));
+          out[k*4+3]=pack4x8unorm(textureLoad(field,p,0));
+        }`;
+  const pipe=d.createComputePipeline({layout:'auto',compute:{module:d.createShaderModule({code}),entryPoint:'main'}});
+  const bytes=R.px.length*16;
+  const storage=d.createBuffer({size:bytes,usage:GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC});
+  const read=d.createBuffer({size:bytes,usage:GPUBufferUsage.COPY_DST|GPUBufferUsage.MAP_READ});
+  const entries=[GPU.colourTex,GPU.backTex,GPU.treeTex,GPU.fieldTex].map((t,binding)=>({binding,resource:t.createView()}));
+  entries.push({binding:4,resource:{buffer:storage}});
+  const bind=d.createBindGroup({layout:pipe.getBindGroupLayout(0),entries});
+  const enc=d.createCommandEncoder(),pass=enc.beginComputePass();pass.setPipeline(pipe);pass.setBindGroup(0,bind);pass.dispatchWorkgroups(Math.ceil(R.px.length/64));pass.end();
+  enc.copyBufferToBuffer(storage,0,read,0,bytes);d.queue.submit([enc.finish()]);
+  await read.mapAsync(GPUMapMode.READ);
+  const actual=new Uint32Array(read.getMappedRange()).slice();
+  read.unmap();storage.destroy();read.destroy();
+  return actual;
+}

test/play.py

@@ -1328,8 +1328,14 @@ with launch(viewport={'width': 1280, 'height': 800}) as (b, pg, url):
     check('fog toggle shows the same rock', fp.evaluate(sample) != 0)
     fp.evaluate("SET.fog = true; SET.look = 'smooth'; SM.t0 = performance.now() + 100000")
     fp.wait_for_timeout(500)
-    check('fog draws in the smooth look', fp.evaluate('R.smoothOn && SM.ok'))
+    check('fog draws in the smooth look', fp.evaluate('R.smoothOn && (R.gpuOn ? GPU.ok : SM.ok)'))
-    check('unseen smooth grains have no terrain shape', fp.evaluate("(() => {const k=(235-R.vy0)*R.bw+300-R.vx0;return R.px[k]===0x04000000 && R.backPx[k]===0 && R.treePx[k]===0 && R.fieldPx[k]===255})()"))
+    fp.add_script_tag(path=str(root / 'test/gpu-readback.js'))
+    check('unseen smooth grains have no terrain shape', fp.evaluate("""async () => {
+      const k=(235-R.vy0)*R.bw+300-R.vx0;
+      const inputs=R.gpuOn ? Array.from((await testSmoothInputs()).slice(k*4,k*4+4)) :
+        [R.px[k],R.backPx[k],R.treePx[k],R.fieldPx[k]];
+      return inputs[0]===0x04000000 && inputs[1]===0 && inputs[2]===0 && inputs[3]===255;
+    }"""))
     fp.screenshot(path=str(out / 'fog-smooth.png'))
     fog_context.close()
 

test/t_gpu.py

@@ -65,10 +65,42 @@ with launch(viewport={'width': 960, 'height': 600}) as (browser, pg, url):
     assert parity['max'] <= 8 and parity['nonblack'] > 1000, parity
     print('PASS world replacement and zoom reallocate textures and draw the new world', parity)
     pg.evaluate('SET.look="smooth"; drawWorld();')
-    assert pg.evaluate('!R.gpuOn && R.smoothOn && GPU.canvas.hidden')
+    assert pg.evaluate('R.gpuOn && R.smoothOn && !GPU.canvas.hidden')
-    print('PASS smooth keeps the legacy path')
+    print('PASS smooth uses WebGPU')
+    pg.evaluate('() => {' + (root / 'tools/scenes/t88-renderer.js').read_text() + '}')
+    pg.add_script_tag(path=str(root / 'test/gpu-readback.js'))
+    # Compare all four smooth inputs, including valid zero tags and class-zero overrides.
+    smooth_inputs = pg.evaluate("""async () => {
+      // Every material also gets a plain, unowned cell (including part materials).
+      for(let m=0;m<NMAT;m++) {
+        const x=CAM[0]-16+(m%11)*3,y=CAM[1]-40+Math.floor(m/11)*2,i=y*W+x;
+        setCell(i,m,20,0);part[i]=mach[i]=occ[i]=bgm[i]=0;blockP[(y>>2)*BW+(x>>2)]=0;
+      }
+      const results=[];
+      const now=performance.now.bind(performance); performance.now=()=>5000; SM.t0=0;
+      try {
+      for(const zoom of [12,6,2])for(const fogOn of [false,true]) {
+        SET.zoom=zoom;SET.fog=fogOn;SET.look='smooth';R.lens=false;SET.lighting='off';SET.glow='off';
+        resizeRender();bot.x=CAM[0]-bot.w/2;bot.y=CAM[1]-bot.h/2;snapCamera();R.frame=999;
+        SET.renderer='legacy';SM.scale=1;drawWorld();
+        const expected=[R.px.slice(),R.backPx.slice(),R.treePx.slice(),R.fieldPx.slice()];
+        SET.renderer='webgpu';drawWorld();
+        const actual=await testSmoothInputs(); let alpha=0,max=0,bad=0,examples=[];
+        for(let k=0;k<R.px.length;k++)for(let t=0;t<4;t++) {
+          const a=expected[t][k],b=actual[k*4+t];if((a>>>24)!==(b>>>24))alpha++;
+          for(let c=0;c<3;c++){const diff=Math.abs((a>>>(c*8)&255)-(b>>>(c*8)&255));max=Math.max(max,diff);if(diff>1){bad++;if(examples.length<3)examples.push({k,t,a,b});}}
+        }
+        results.push({zoom,fogOn,alpha,max,bad,examples,share:GPU.overrideCount/R.px.length});
+      }
+      return results;
+      } finally { performance.now=now; }
+    }""")
+    for r in smooth_inputs:
+        assert r['alpha'] == 0 and r['bad'] == 0, r
+    print('PASS all four smooth inputs match legacy: classes and flags exact, RGB within 1', smooth_inputs)
     pg.evaluate('SET.look="pixel"; SET.renderer="webgpu"; PERF_MARKS=true; drawWorld();')
-    pg.wait_for_function('!GPU.timingBusy')
+    # This fixture stops requestAnimationFrame; poll asynchronous GPU events by timer.
+    pg.wait_for_function('!GPU.timingBusy', polling=10)
     assert pg.evaluate('!GPU.queries || GPU.gpuMs>=0')
     assert pg.evaluate('perfMarks.override>=0 && perfMarks.upload>=0')
     print('PASS separate override/upload clocks and optional GPU timestamps')
@@ -103,9 +135,23 @@ with launch(viewport={'width': 960, 'height': 600}) as (browser, pg, url):
     }""")
     assert all(cadence.values()), cadence
     print('PASS GPU-only lighting, both resolutions and lightEvery cadence')
+    switches = pg.evaluate("""() => {
+      const original=computeLight,frames=[];
+      computeLight=()=>{throw Error('CPU light called in WebGPU smooth');};
+      try {
+        SET.lightEvery=100;
+        for(const look of ['smooth','pixel','smooth']) {
+          SET.look=look;R.frame++;drawWorld();
+          frames.push(R.gpuOn && GPU.lightFrame===R.frame && R.smoothOn===(look==='smooth'));
+        }
+        return frames;
+      } finally {computeLight=original;SET.lightEvery=1;}
+    }""")
+    assert all(switches), switches
+    print('PASS switching either look refreshes new light textures without CPU lighting')
     assert pg.evaluate('!GPU.queries || GPU.lightMs>=0')
     pg.evaluate('GPU.device.destroy()')
-    pg.wait_for_function('!GPU.ok')
+    pg.wait_for_function('!GPU.ok', polling=10)
     assert pg.evaluate('SET.renderer==="legacy"')
     pg.evaluate('drawWorld()')
     print('PASS device loss restores the legacy renderer')

test/t_smooth.py

@@ -1,4 +1,4 @@
-"""The smooth look: the WebGL pass compiles, draws, falls back cleanly, and the thermal lens still works."""
+"""The smooth look: the WebGPU pass compiles, draws, falls back cleanly, and the thermal lens still works."""
 import sys, pathlib
 sys.path.insert(0, str(pathlib.Path(__file__).resolve().parent.parent / 'tools'))
 from use_venv import ensure_playwright; ensure_playwright()
@@ -15,11 +15,11 @@ with launch(viewport={'width': 1024, 'height': 640}) as (b, pg, url):
     pg.on('pageerror', lambda e: logs.append(f'PAGEERROR: {e}'))
     pg.on('console', lambda m: logs.append(f'{m.type}: {m.text}') if m.type in ('error', 'warning') else None)
     pg.add_init_script('localStorage.clear();')
-    pg.goto(url); pg.evaluate('SET.fog = false'); pg.wait_for_timeout(300)
+    pg.goto(url); pg.evaluate('CAPS_CHECK'); pg.evaluate('SET.fog = false'); pg.wait_for_timeout(300)
     pg.click('#btn-new'); pg.fill('#nw-seed', '777'); pg.click('#btn-generate')
     pg.wait_for_function('GAME.running === true', timeout=20000); pg.wait_for_timeout(1500)
     ev = pg.evaluate
-    check('the smooth look is the default and is drawing', ev("SET.look === 'smooth' && SM.ok && R.smoothOn"), ev('[SET.look, SM.ok, SM.why, R.smoothOn]'))
+    check('the smooth look is the default and is drawing', ev("SET.look === 'smooth' && GPU.ok && R.gpuOn && R.smoothOn"), ev('[SET.look, SM.ok, SM.why, R.smoothOn]'))
     # the canvas shows the GPU image: sample a sky pixel and a ground pixel, they must differ and not be black
     px = ev("""(() => { const c = worldPixelContext().getImageData(R.canvas.width >> 1, 4, 1, 1).data, g = worldToScreenPx(bot.x + 2, surfY[Math.floor(bot.x)] + 6), d = worldPixelContext().getImageData(Math.round(g.x), Math.round(g.y), 1, 1).data; return [Array.from(c), Array.from(d)]; })()""")
     check('it draws sky and ground', sum(px[0][:3]) > 200 and px[0][:3] != px[1][:3], px)
@@ -32,8 +32,10 @@ with launch(viewport={'width': 1024, 'height': 640}) as (b, pg, url):
     check('the pixel look can be chosen', ev('!R.smoothOn'))
     pg.screenshot(path=str(out / 'pixel.png'))
     ev("SET.look = 'smooth'"); pg.wait_for_timeout(1500)
-    ev("SM.scale = 0.5"); pg.wait_for_timeout(1500)
+    ev("SET.renderer = 'legacy'; SM.scale = 0.5"); pg.wait_for_timeout(1500)
     st = ev('[R.smoothOn, SM.scale, SM.canvas.width, R.canvas.width, SET.look, SM.ok, SM.why]'); check('half-resolution smooth still draws', (st[0] and st[2] == round(st[3] * st[1])) or st[1] > 0.5, st)
+    ev("SET.renderer = 'webgpu'"); pg.wait_for_timeout(300)
+    check('smooth returns to WebGPU after legacy comparison', ev('R.gpuOn && R.smoothOn'))
     # Uniform colours isolate the solid/loose seam from material noise and lighting.
     ground = ev('''() => {
       GAME.paused = true; requestAnimationFrame = () => 0;
@@ -290,7 +292,7 @@ with launch(viewport={'width': 1024, 'height': 640}) as (b, pg, url):
 # a browser without hardware drawing gets the pixel look and a note
 with launch(gpu='soft', viewport={'width': 800, 'height': 500}) as (_, pg2, soft_url):
     pg2.add_init_script('window.__allowSoftwareGL = false; localStorage.clear();')
-    pg2.goto(soft_url); pg2.evaluate('SET.fog = false'); pg2.wait_for_timeout(300)
+    pg2.goto(soft_url); pg2.evaluate("SET.fog = false; SET.renderer='legacy'"); pg2.wait_for_timeout(300)
     pg2.click('#btn-new'); pg2.fill('#nw-seed', '5'); pg2.click('#btn-generate')
     pg2.wait_for_function('GAME.running === true', timeout=20000); pg2.wait_for_timeout(1200)
     check('software rendering falls back to pixels with a note', pg2.evaluate("SET.look === 'pixel' && !R.smoothOn && SM.why === 'software'") and pg2.locator('.toast', has_text='graphics card').count() >= 1, pg2.evaluate('[SET.look, SM.why]'))

tools/perf-measure.js

@@ -66,7 +66,7 @@ async function perfRender(lighting, overlays) {
     return { frame_ms, drawWorld_ms: draw, computeLight_ms: light, marks_ms: marks,
       overlays: Object.fromEntries(Object.entries(data).map(([key, d]) => [key, { ms_per_frame: d.ms / n, calls: d.calls }])),
       overlay_ms: Object.values(data).reduce((sum, d) => sum + d.ms / n, 0),
-      scale: SM.scale, actual_look: SET.look, renderer: R.gpuOn ? 'webgpu' : 'legacy', override_cells: GPU.overrideCount || 0, gpu_ms };
+      scale: R.gpuOn ? 1 : SM.scale, actual_look: SET.look, renderer: R.gpuOn ? 'webgpu' : 'legacy', override_cells: GPU.overrideCount || 0, gpu_ms };
   } finally { PERF_MARKS = false; }
 }
 function perfRun(seconds) {
@@ -84,7 +84,7 @@ function perfRun(seconds) {
     requestAnimationFrame(frame);
     setTimeout(() => {
       window.__perfHalt(); PERF_FRAME = null; GAME.paused = true;
-      resolve({ samples, end: perfState(), scale: SM.scale, actual_look: SET.look, renderer: R.gpuOn ? 'webgpu' : 'legacy', override_cells: GPU.overrideCount || 0, gpu_ms: R.gpuOn ? GPU.gpuMs : null });
+      resolve({ samples, end: perfState(), scale: R.gpuOn ? 1 : SM.scale, actual_look: SET.look, renderer: R.gpuOn ? 'webgpu' : 'legacy', override_cells: GPU.overrideCount || 0, gpu_ms: R.gpuOn ? GPU.gpuMs : null });
     }, seconds * 1000);
   });
 }

tools/perf-scenes.js

@@ -61,6 +61,6 @@ function perfView(look, zoom) {
   SET.look = look; SET.zoom = zoom; SET.lighting = 'full';
   resizeRender(); snapCamera();
   for (let k = 0; k < 5; k++) { R.frame++; drawWorld(); }
-  if (look === 'smooth' && (!SM.ok || !R.smoothOn)) throw Error('Smooth look unavailable: ' + SM.why);
+  if (look === 'smooth' && (!R.smoothOn || (!R.gpuOn && !SM.ok))) throw Error('Smooth look unavailable: ' + SM.why);
-  return { grains: R.bw * R.bh, buffer: [R.bw, R.bh], scale: SM.scale };
+  return { grains: R.bw * R.bh, buffer: [R.bw, R.bh], scale: R.gpuOn ? 1 : SM.scale };
 }

tools/perf.py

@@ -191,7 +191,7 @@ def main():
                         view['lighting'][light] = pg.evaluate('([l,o]) => perfRender(l,o)', [light, overlays])
                     if view['lighting']['full']['actual_look'] != look or view['lighting']['off']['actual_look'] != look:
                         raise RuntimeError('Look fell back while measuring')
-                    if look == 'pixel' and any(r['renderer'] != 'webgpu' for r in view['lighting'].values()):
+                    if any(r['renderer'] != 'webgpu' for r in view['lighting'].values()):
                         raise RuntimeError('WebGPU fell back while measuring: ' + pg.evaluate('GPU.error'))
                     data['views'][f'{look}/{zoom:g}'] = view
             for look in looks:
@@ -201,7 +201,7 @@ def main():
                     r = pg.evaluate('seconds => perfRun(seconds)', args.seconds)
                     if r['actual_look'] != look:
                         raise RuntimeError('Look fell back during real run')
-                    if look == 'pixel' and r['renderer'] != 'webgpu':
+                    if r['renderer'] != 'webgpu':
                         raise RuntimeError('WebGPU fell back during real run: ' + pg.evaluate('GPU.error'))
                     for key in ['interval', 'work', 'sim']:
                         r[key + '_ms'] = quantiles([s[key] for s in r['samples']])

tools/scenes/t88-renderer.js

@@ -11,13 +11,15 @@ for (let y = cy - 52; y < cy + 52; y++) for (let x = cx - 90; x < cx + 90; x++)
 for (let x = cx - 90; x < cx + 90; x++) { surfY[x] = cy + 16; dirtY[x] = cy + 30; graniteY[x] = cy + 38; basaltY[x] = cy + 46; }
 const fill = (x0,y0,w,h,m,T=20) => { for(let y=y0;y<y0+h;y++)for(let x=x0;x<x0+w;x++)setCell(y*W+x,m,T,F_NAT); };
 fill(cx-62,cy+6,22,16,WATER); fill(cx-34,cy+11,16,11,MAGMA,1450);
-fill(cx-33,cy+3,8,8,FIRE,900); fill(cx-12,cy+17,15,6,SAND);
+fill(cx-33,cy+3,8,8,FIRE,900); fill(cx-32,cy-4,7,6,SMOKE); fill(cx-12,cy+17,15,6,SAND);
 fill(cx-10,cy+24,12,5,STONE,1100);
 for(let y=cy-20;y<cy+16;y++)for(let x=cx-58;x<cx-54;x++)bgm[y*W+x]=WOOD;
 for(let y=cy-28;y<cy-12;y++)for(let x=cx-68;x<cx-44;x++)if((x-cx+56)**2+(y-cy+20)**2<130)bgm[y*W+x]=LEAVES;
 for(let x=cx-40;x<cx-34;x++)digp[(cy+24)*W+x]=160;
 const place=(code,v,dx,dy)=>{const e=placeBlock(code,v,(cx+dx)>>2,(cy+dy)>>2,true);if(e)throw Error(e);};
 for(let dx=-36;dx<-12;dx+=4)place(BP_BELT,1,dx,-12);
+place(BP_FILTER,128,-12,-12); filterCfg.set(((cy-12)>>2)*BW+((cx-12)>>2),{m:[SAND],inv:1});
+place(BP_SLOPE_BL,RSTONE,44,-12);
 place(BP_SCREW,1,-4,-12); place(BP_SCREW,4,0,-12);
 place(BP_BELT_UP,1,16,-12);place(BP_BELT_UP,1,20,-16);
 place(BP_PIPE,0,32,-12);place(BP_PIPE,0,36,-12);

tools/shot.py

@@ -92,6 +92,8 @@ def main():
         for code in [pathlib.Path(a.setup).read_text() if a.setup else None, a.js]:
             if code:
                 pg.evaluate(f'() => {{ {code}\n}}')
+        # Setup scenes can change surfY after the renderer cached the world's horizon.
+        pg.evaluate('SM.surf = -1')
         if a.ticks:
             pg.evaluate(f'for (let k = 0; k < {a.ticks}; k++) simTick();')
         pg.evaluate(f"""GAME.paused = true; $('hud').hidden = {'false' if a.hud else 'true'}; GAME.hidePreviews = {'false' if a.previews else 'true'};
@@ -103,7 +105,7 @@ def main():
             pg.evaluate("() => { window.__frozenNow = 5000; window.requestAnimationFrame = () => 0; }")
             pg.wait_for_timeout(100)
             pg.evaluate("() => { R.frame = 999; if (typeof SM !== 'undefined') SM.t0 = 0; frame(lastRender + 1000); }")
-        if a.look == 'pixel' and a.renderer == 'webgpu' and not pg.evaluate('R.gpuOn'):
+        if a.renderer == 'webgpu' and not pg.evaluate('R.gpuOn'):
             raise RuntimeError('WebGPU fell back: ' + pg.evaluate('GPU.error'))
         kw = {}
         if a.clip:
@@ -111,9 +113,9 @@ def main():
             kw['clip'] = {'x': x, 'y': y, 'width': w, 'height': h}
         pathlib.Path(a.out).parent.mkdir(parents=True, exist_ok=True)
         pg.screenshot(path=a.out, **kw)
-        note = pg.evaluate("SET.look === 'smooth' ? (SM.ok ? 'smooth look' : 'smooth look FAILED: ' + SM.why) : 'pixel look'")
+        note = pg.evaluate("SET.look === 'smooth' ? (R.smoothOn ? 'smooth look' : 'smooth look FAILED: ' + SM.why) : 'pixel look'")
         b.close()
-    print(f"saved {a.out} ({note}, {a.renderer if a.look == 'pixel' else 'legacy'})")
+    print(f"saved {a.out} ({note}, {a.renderer})")
     for e in errors:
         print('ERROR', e)
     sys.exit(1 if errors else 0)