26 changed files with 336 additions and 2077 deletions
--- a/additional/dom-webgpu.js
+++ b/additional/dom-webgpu.js
@ -1,885 +0,0 @@
 /*
 Crafter.Graphics WebGPU bridge — DOM mode UI compute pipeline.
 Surface model (high-level, deliberately not 1:1 with WebGPU):
  - JS owns the GPUDevice/queue/compute pipelines/bind-group cache.
  - C++ owns slot allocation and per-frame logic; it calls into ~15 imports.
  - Standard UI shaders are embedded as WGSL strings at the bottom of this
    file and compiled once at startup.
 Ping-pong output strategy (per Decision 2 in plan):
  - Two rgba8unorm storage textures sized to the canvas.
  - Each Dispatch reads `prev` (sampled), writes `out` (storage, write-only).
  - JS swaps the two between dispatches, so item-order overdraw works.
  - At frame end, the current `out` is blitted to the canvas via
    copyTextureToTexture (canvas configured as rgba8unorm to match).
 This file is loaded as <script type="module">. Top-level await blocks
 runtime.js's _start() until adapter + device are resolved, so by the time
 main() runs every import here is fully wired.
 */
 // ─── env stubs (assigned synchronously, BEFORE any async work) ────────────
 // The wasm module's import-link step needs every declared wgpu* import to
 // resolve to a Function. If init below throws, the stubs stay in place so
 // the wasm still links — and the call site gets a clear error at runtime
 // instead of "import object field X is not a Function" at link time.
 window.crafter_webbuild_env = window.crafter_webbuild_env || {};
 window.crafter_webbuild_env.table = window.crafter_webbuild_env.table
    || new WebAssembly.Table({ initial: 4, element: "anyfunc" });
 let initError = null;
 function stub(name) {
    return (...args) => {
        const msg = `[crafter-wgpu] ${name}() called but WebGPU init failed: ${initError?.message ?? "(no error captured)"}`;
        console.error(msg);
        throw new Error(msg);
    };
 }
 {
    const e = window.crafter_webbuild_env;
    for (const n of [
        "wgpuGetCanvasWidth", "wgpuGetCanvasHeight", "wgpuSurfaceWidth", "wgpuSurfaceHeight",
        "wgpuInit", "wgpuCreateBuffer", "wgpuWriteBuffer", "wgpuDestroyBuffer",
        "wgpuCreateAtlasTexture", "wgpuWriteAtlasRegion", "wgpuDestroyTexture",
        "wgpuCreateLinearClampSampler", "wgpuFrameBegin", "wgpuFrameEnd",
        "wgpuDispatchQuads", "wgpuDispatchCircles", "wgpuDispatchImages", "wgpuDispatchText",
    ]) {
        // Read-write ints don't need a stub-throw; return 0 for the size queries.
        e[n] = n.endsWith("Width") || n.endsWith("Height") ? () => 0 : stub(n);
    }
 }
 // ─── canvas + device init (runs before _start) ───────────────────────────
 // Wrapped in an async IIFE assigned to window.crafter_webbuild_env_ready so
 // the runtime.js shim can `await` it explicitly before calling _start().
 // Sibling <script type="module"> top-level awaits are NOT reliably
 // serialized in Firefox (verified 2026-05), so we can't depend on this
 // file's TLA to block runtime.js by itself.
 window.crafter_webbuild_env_ready = (async () => {
 try {
 if (!navigator.gpu) {
    document.body.innerHTML = "<p style=\"font-family:sans-serif;padding:24px\">"
        + "WebGPU not available in this browser. Try Chrome 121+ / Firefox 141+ / Safari 26+.</p>";
    initError = new Error("WebGPU unavailable");
    throw initError;
 }
 const canvas = document.createElement("canvas");
 canvas.id = "crafter-canvas";
 canvas.style.cssText = "position:fixed;inset:0;width:100vw;height:100vh;display:block;";
 document.body.style.margin = "0";
 document.body.appendChild(canvas);
 function syncCanvasSize() {
    // Match canvas pixel size to its CSS pixel size 1:1 so MouseEvent
    // clientX/clientY (CSS pixels) and the wasm-side window.width/.height
    // share the same coordinate space. (HiDPI sharpness is a v2 concern
    // — would need DPR on the GPU side AND a scaling step in the C++
    // Window/Event glue so layout/hit-testing/dispatch counts stay
    // consistent.)
    const w = window.innerWidth;
    const h = window.innerHeight;
    if (canvas.width !== w)  canvas.width  = w;
    if (canvas.height !== h) canvas.height = h;
    return { w, h };
 }
 syncCanvasSize();
 const adapter = await navigator.gpu.requestAdapter();
 if (!adapter) {
    initError = new Error("navigator.gpu.requestAdapter() returned null (no compatible adapter)");
    console.error("[crafter-wgpu]", initError.message);
    throw initError;
 }
 const device  = await adapter.requestDevice();
 const queue   = device.queue;
 const ctx     = canvas.getContext("webgpu");
 const canvasFormat = "rgba8unorm"; // match storage textures, skip swizzle blit
 ctx.configure({ device, format: canvasFormat, alphaMode: "opaque",
                usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_DST });
 device.lost.then((info) => {
    console.error("[crafter-wgpu] device lost:", info.message);
    state.gpuLost = true;
 });
 // ─── handle tables ─────────────────────────────────────────────────────
 const buffers       = new Map();   // handle → GPUBuffer
 const textures      = new Map();   // handle → GPUTexture
 const textureViews  = new Map();   // handle → GPUTextureView  (mirrors textures key for the view)
 const samplers      = new Map();   // handle → GPUSampler
 let nextHandle = 1;
 function newHandle() { return nextHandle++; }
 // ─── ping-pong storage textures ────────────────────────────────────────
 const state = {
    pingTex: null, pingView: null,
    pongTex: null, pongView: null,
    outIsPing: true,  // current "out" target
    width: 0, height: 0,
    encoder: null,
    pass: null,
    headerRing: null,        // GPUBuffer; uniform header writes ring through this
    headerRingSize: 0,
    headerRingOffset: 0,
    bindGroupCache: new Map(),  // key → GPUBindGroup
    gpuLost: false,
 };
 function recreatePingPong(w, h) {
    const usage = GPUTextureUsage.STORAGE_BINDING
                | GPUTextureUsage.TEXTURE_BINDING
                | GPUTextureUsage.COPY_SRC
                | GPUTextureUsage.COPY_DST;   // COPY_DST so we can clear it
    if (state.pingTex) state.pingTex.destroy();
    if (state.pongTex) state.pongTex.destroy();
    state.pingTex = device.createTexture({ size: [w, h], format: "rgba8unorm", usage });
    state.pongTex = device.createTexture({ size: [w, h], format: "rgba8unorm", usage });
    state.pingView = state.pingTex.createView();
    state.pongView = state.pongTex.createView();
    state.width = w; state.height = h;
    state.outIsPing = true;
    state.bindGroupCache.clear();
 }
 function ensureSized() {
    const { w, h } = syncCanvasSize();
    if (w !== state.width || h !== state.height) {
        recreatePingPong(w, h);
        // Notify the wasm side that the surface size changed so it can
        // fire onResize through Window. The wasm export is added by
        // Crafter.Graphics-Window.cpp.
        const onResize = wasmExports && wasmExports.__crafterDom_resize;
        if (onResize) onResize(1, w, h);
    }
 }
 // Header ring buffer: 256-byte-aligned slots holding UIDispatchHeader (48
 // bytes of meaningful data, padded to 256). Wraps at frame boundary.
 const HEADER_ALIGN = 256;
 const HEADER_RING_SLOTS = 64;
 state.headerRingSize = HEADER_ALIGN * HEADER_RING_SLOTS;
 state.headerRing = device.createBuffer({
    size: state.headerRingSize,
    usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
 });
 // ─── pipelines ─────────────────────────────────────────────────────────
 const wgslShared = String.raw`
 struct UIDispatchHeader {
    outImage:      u32,
    itemBuffer:    u32,
    surfaceW:      u32,
    surfaceH:      u32,
    clipX:         f32,
    clipY:         f32,
    clipW:         f32,
    clipH:         f32,
    itemCount:     u32,
    frameIdx:      u32,
    flags:         u32,
    _pad:          u32,
 };
@group(0) @binding(0) var<uniform> hdr : UIDispatchHeader;
@group(1) @binding(0) var outTex  : texture_storage_2d<rgba8unorm, write>;
@group(1) @binding(1) var prevTex : texture_2d<f32>;
 fn uiResolvePixel(coord: vec2<u32>) -> bool {
    if (coord.x >= hdr.surfaceW || coord.y >= hdr.surfaceH) { return false; }
    let fx = f32(coord.x); let fy = f32(coord.y);
    if (fx < hdr.clipX || fy < hdr.clipY) { return false; }
    if (fx >= hdr.clipX + hdr.clipW) { return false; }
    if (fy >= hdr.clipY + hdr.clipH) { return false; }
    return true;
 }
 fn uiBlendOver(dst: vec4<f32>, src: vec4<f32>) -> vec4<f32> {
    let a = clamp(src.a, 0.0, 1.0);
    let rgb = mix(dst.rgb, src.rgb, vec3<f32>(a));
    let outA = a + dst.a * (1.0 - a);
    return vec4<f32>(rgb, outA);
 }
 fn uiSdRoundRect(p: vec2<f32>, halfSize: vec2<f32>, r4: vec4<f32>) -> f32 {
    var r: vec4<f32> = r4;
    // Pick radius for the quadrant p is in. r order: (TL, TR, BR, BL).
    let rx = select(r.x, r.z, p.x > 0.0);
    let ry = select(r.w, r.y, p.x > 0.0);
    let radius = select(ry, rx, p.y > 0.0);
    let q = abs(p) - halfSize + vec2<f32>(radius);
    return min(max(q.x, q.y), 0.0) + length(max(q, vec2<f32>(0.0))) - radius;
 }
 `;
 const wgslQuads = wgslShared + String.raw`
 struct QuadItem {
    rect:    vec4<f32>,
    color:   vec4<f32>,
    corners: vec4<f32>,
    outline: vec4<f32>,
 };
@group(2) @binding(0) var<storage, read> items : array<QuadItem>;
@compute @workgroup_size(8, 8, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
    if (!uiResolvePixel(gid.xy)) { return; }
    let coord = vec2<i32>(i32(gid.x), i32(gid.y));
    var dst = textureLoad(prevTex, coord, 0);
    let sp = vec2<f32>(f32(gid.x) + 0.5, f32(gid.y) + 0.5);
    for (var i: u32 = 0u; i < hdr.itemCount; i = i + 1u) {
        let it = items[i];
        let lo = it.rect.xy;
        let hi = it.rect.xy + it.rect.zw;
        if (sp.x < lo.x || sp.y < lo.y || sp.x >= hi.x || sp.y >= hi.y) { continue; }
        let halfSize = it.rect.zw * 0.5;
        let p = sp - (it.rect.xy + halfSize);
        let d = uiSdRoundRect(p, halfSize, it.corners);
        let bodyA = clamp(0.5 - d, 0.0, 1.0);
        if (bodyA <= 0.0 && it.outline.x <= 0.0) { continue; }
        var src = vec4<f32>(it.color.rgb, it.color.a * bodyA);
        if (it.outline.x > 0.0) {
            let t = abs(d + it.outline.x * 0.5) - it.outline.x * 0.5;
            let outlineA = clamp(0.5 - t, 0.0, 1.0);
            src = vec4<f32>(mix(src.rgb, it.outline.yzw, vec3<f32>(outlineA)),
                            max(src.a, outlineA));
        }
        if (src.a <= 0.0) { continue; }
        dst = uiBlendOver(dst, src);
    }
    textureStore(outTex, coord, dst);
 }
 `;
 const wgslCircles = wgslShared + String.raw`
 struct CircleItem {
    centerRadius: vec4<f32>,
    color:        vec4<f32>,
    outline:      vec4<f32>,
 };
@group(2) @binding(0) var<storage, read> items : array<CircleItem>;
@compute @workgroup_size(8, 8, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
    if (!uiResolvePixel(gid.xy)) { return; }
    let coord = vec2<i32>(i32(gid.x), i32(gid.y));
    var dst = textureLoad(prevTex, coord, 0);
    let sp = vec2<f32>(f32(gid.x) + 0.5, f32(gid.y) + 0.5);
    for (var i: u32 = 0u; i < hdr.itemCount; i = i + 1u) {
        let it = items[i];
        let center = it.centerRadius.xy;
        let radius = it.centerRadius.z;
        let d = length(sp - center) - radius;
        let bodyA = clamp(0.5 - d, 0.0, 1.0);
        if (bodyA <= 0.0 && it.outline.x <= 0.0) { continue; }
        var src = vec4<f32>(it.color.rgb, it.color.a * bodyA);
        if (it.outline.x > 0.0) {
            let t = abs(d + it.outline.x * 0.5) - it.outline.x * 0.5;
            let outlineA = clamp(0.5 - t, 0.0, 1.0);
            src = vec4<f32>(mix(src.rgb, it.outline.yzw, vec3<f32>(outlineA)),
                            max(src.a, outlineA));
        }
        if (src.a <= 0.0) { continue; }
        dst = uiBlendOver(dst, src);
    }
    textureStore(outTex, coord, dst);
 }
 `;
 const wgslImages = wgslShared + String.raw`
 struct ImageItem {
    rect:  vec4<f32>,
    uv:    vec4<f32>,
    tint:  vec4<f32>,
    slots: vec4<u32>,
 };
@group(2) @binding(0) var<storage, read> items : array<ImageItem>;
@group(3) @binding(0) var imgTex : texture_2d<f32>;
@group(3) @binding(1) var imgSampler : sampler;
@compute @workgroup_size(8, 8, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
    if (!uiResolvePixel(gid.xy)) { return; }
    let coord = vec2<i32>(i32(gid.x), i32(gid.y));
    var dst = textureLoad(prevTex, coord, 0);
    let sp = vec2<f32>(f32(gid.x) + 0.5, f32(gid.y) + 0.5);
    for (var i: u32 = 0u; i < hdr.itemCount; i = i + 1u) {
        let it = items[i];
        let lo = it.rect.xy;
        let hi = it.rect.xy + it.rect.zw;
        if (sp.x < lo.x || sp.y < lo.y || sp.x >= hi.x || sp.y >= hi.y) { continue; }
        let t = (sp - lo) / it.rect.zw;
        let uv = vec2<f32>(mix(it.uv.x, it.uv.z, t.x), mix(it.uv.y, it.uv.w, t.y));
        let sample = textureSampleLevel(imgTex, imgSampler, uv, 0.0);
        let src = sample * it.tint;
        if (src.a <= 0.0) { continue; }
        dst = uiBlendOver(dst, src);
    }
    textureStore(outTex, coord, dst);
 }
 `;
 const wgslText = wgslShared + String.raw`
 struct GlyphItem {
    rect:  vec4<f32>,
    uv:    vec4<f32>,
    color: vec4<f32>,
 };
@group(2) @binding(0) var<storage, read> items : array<GlyphItem>;
@group(3) @binding(0) var atlasTex : texture_2d<f32>;
@group(3) @binding(1) var atlasSampler : sampler;
@compute @workgroup_size(8, 8, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
    if (!uiResolvePixel(gid.xy)) { return; }
    let coord = vec2<i32>(i32(gid.x), i32(gid.y));
    var dst = textureLoad(prevTex, coord, 0);
    let sp = vec2<f32>(f32(gid.x) + 0.5, f32(gid.y) + 0.5);
    for (var i: u32 = 0u; i < hdr.itemCount; i = i + 1u) {
        let it = items[i];
        let lo = it.rect.xy;
        let hi = it.rect.xy + it.rect.zw;
        if (sp.x < lo.x || sp.y < lo.y || sp.x >= hi.x || sp.y >= hi.y) { continue; }
        let t = (sp - lo) / it.rect.zw;
        let uv = vec2<f32>(mix(it.uv.x, it.uv.z, t.x), mix(it.uv.y, it.uv.w, t.y));
        // stb_truetype SDF: pixel value ~128 is the edge. Treat alpha as
        // the smoothed step around that midpoint.
        let sdf = textureSampleLevel(atlasTex, atlasSampler, uv, 0.0).r;
        let alpha = clamp((sdf - 0.5) * 8.0 + 0.5, 0.0, 1.0);
        if (alpha <= 0.0) { continue; }
        let src = vec4<f32>(it.color.rgb, it.color.a * alpha);
        dst = uiBlendOver(dst, src);
    }
    textureStore(outTex, coord, dst);
 }
 `;
 function makePipeline(label, wgsl, hasGroup3) {
    const mod = device.createShaderModule({ label, code: wgsl });
    // Layout: group 0 uniform header, group 1 (out storage + prev sampled),
    // group 2 storage items SSBO, optional group 3 (texture + sampler).
    const bgl0 = device.createBindGroupLayout({ entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: "uniform", hasDynamicOffset: true, minBindingSize: 48 } },
    ]});
    const bgl1 = device.createBindGroupLayout({ entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, storageTexture: { format: "rgba8unorm", access: "write-only", viewDimension: "2d" } },
        { binding: 1, visibility: GPUShaderStage.COMPUTE, texture: { sampleType: "float", viewDimension: "2d" } },
    ]});
    const bgl2 = device.createBindGroupLayout({ entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: "read-only-storage" } },
    ]});
    const layouts = [bgl0, bgl1, bgl2];
    let bgl3 = null;
    if (hasGroup3) {
        bgl3 = device.createBindGroupLayout({ entries: [
            { binding: 0, visibility: GPUShaderStage.COMPUTE, texture: { sampleType: "float", viewDimension: "2d" } },
            { binding: 1, visibility: GPUShaderStage.COMPUTE, sampler: { type: "filtering" } },
        ]});
        layouts.push(bgl3);
    }
    const pl = device.createPipelineLayout({ bindGroupLayouts: layouts });
    const pipeline = device.createComputePipeline({
        layout: pl,
        compute: { module: mod, entryPoint: "main" },
    });
    return { pipeline, bgl0, bgl1, bgl2, bgl3 };
 }
 const pipeQuads   = makePipeline("ui-quads",   wgslQuads,   false);
 const pipeCircles = makePipeline("ui-circles", wgslCircles, false);
 const pipeImages  = makePipeline("ui-images",  wgslImages,  true);
 const pipeText    = makePipeline("ui-text",    wgslText,    true);
 // Bind groups for group 0 (header uniform with dynamic offset) — one per
 // pipeline, references the same ring buffer.
 function makeHdrBG(pipe) {
    return device.createBindGroup({
        layout: pipe.bgl0,
        entries: [{ binding: 0, resource: { buffer: state.headerRing, offset: 0, size: 48 } }],
    });
 }
 const hdrBG = {
    quads:   makeHdrBG(pipeQuads),
    circles: makeHdrBG(pipeCircles),
    images:  makeHdrBG(pipeImages),
    text:    makeHdrBG(pipeText),
 };
 // Group 1 changes between dispatches because `out` and `prev` swap on the
 // ping-pong. Cached by current ping-pong direction and texture size; the
 // stored bind group is reusable across all pipelines that share a
 // layout-compatible bgl1 (all standard pipelines and custom shaders do,
 // since they declare identical group-1 entries per the contract).
 function getGroup1BG(bgl1) {
    const key = `g1/${state.outIsPing ? 1 : 0}/${state.width}x${state.height}`;
    let bg = state.bindGroupCache.get(key);
    if (bg) return bg;
    const outView  = state.outIsPing ? state.pingView : state.pongView;
    const prevView = state.outIsPing ? state.pongView : state.pingView;
    bg = device.createBindGroup({
        layout: bgl1,
        entries: [
            { binding: 0, resource: outView },
            { binding: 1, resource: prevView },
        ],
    });
    state.bindGroupCache.set(key, bg);
    return bg;
 }
 function getGroup2BG(pipe, itemsHandle) {
    const key = `items/${pipe === pipeQuads ? "q" : pipe === pipeCircles ? "c" : pipe === pipeImages ? "i" : "t"}/${itemsHandle}`;
    let bg = state.bindGroupCache.get(key);
    if (bg) return bg;
    const buf = buffers.get(itemsHandle);
    if (!buf) throw new Error(`getGroup2BG: unknown items buffer ${itemsHandle}`);
    bg = device.createBindGroup({
        layout: pipe.bgl2,
        entries: [{ binding: 0, resource: { buffer: buf } }],
    });
    state.bindGroupCache.set(key, bg);
    return bg;
 }
 function getGroup3BG(pipe, texHandle, sampHandle) {
    const key = `t3/${texHandle}/${sampHandle}/${pipe === pipeImages ? "i" : "x"}`;
    let bg = state.bindGroupCache.get(key);
    if (bg) return bg;
    const tex = textureViews.get(texHandle);
    const sam = samplers.get(sampHandle);
    if (!tex || !sam) throw new Error(`getGroup3BG: unknown view ${texHandle} / sampler ${sampHandle}`);
    bg = device.createBindGroup({
        layout: pipe.bgl3,
        entries: [
            { binding: 0, resource: tex },
            { binding: 1, resource: sam },
        ],
    });
    state.bindGroupCache.set(key, bg);
    return bg;
 }
 // ─── wasm import surface ───────────────────────────────────────────────
 let wasmExports = null;
 // Crafter.Build's runtime.js exposes the wasi instance on
 // window.crafter_wasi after instantiation. We grab the exports lazily so
 // every import-side function works regardless of call order. memU8 /
 // memF32 / memU32 always re-derive the typed-array view because the
 // wasm memory's backing ArrayBuffer is detached and replaced whenever
 // the wasm grows its memory; caching a typed array would alias to
 // freed memory after a grow.
 function getExports() {
    if (wasmExports) return wasmExports;
    const wasi = window.crafter_wasi;
    if (!wasi || !wasi.instance) {
        throw new Error("[crafter-wgpu] wasm exports not available yet (called too early)");
    }
    wasmExports = wasi.instance.exports;
    return wasmExports;
 }
 function memU8()  { return new Uint8Array(getExports().memory.buffer); }
 function memF32() { return new Float32Array(getExports().memory.buffer); }
 function memU32() { return new Uint32Array(getExports().memory.buffer); }
 // Stubs were assigned at the top of this file; we now overwrite them with
 // real implementations now that init has succeeded.
 const env = window.crafter_webbuild_env;
 env.wgpuGetCanvasWidth  = () => canvas.width;
 env.wgpuGetCanvasHeight = () => canvas.height;
 env.wgpuCreateBuffer = (byteSize) => {
    const h = newHandle();
    const buf = device.createBuffer({
        size: Math.max(16, byteSize),
        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
    });
    buffers.set(h, buf);
    return h;
 };
 env.wgpuWriteBuffer = (handle, srcPtr, byteSize) => {
    state.writeBufferCount = (state.writeBufferCount || 0) + 1;
    state.lastWriteHandle = handle;
    state.lastWriteSize = byteSize;
    const buf = buffers.get(handle);
    if (!buf) return;
    // writeBuffer requires a multiple of 4 bytes and an aligned offset.
    const aligned = (byteSize + 3) & ~3;
    queue.writeBuffer(buf, 0, memU8().buffer, srcPtr, aligned);
 };
 env.wgpuDestroyBuffer = (handle) => {
    const buf = buffers.get(handle);
    if (buf) { buf.destroy(); buffers.delete(handle); }
    // Invalidate any cached bind group that referenced this handle.
    for (const k of state.bindGroupCache.keys()) {
        if (k.startsWith("items/") && k.endsWith("/" + handle)) {
            state.bindGroupCache.delete(k);
        }
    }
 };
 env.wgpuCreateAtlasTexture = (w, h) => {
    const handle = newHandle();
    const tex = device.createTexture({
        size: [w, h],
        format: "r8unorm",
        usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_DST,
    });
    textures.set(handle, tex);
    textureViews.set(handle, tex.createView());
    return handle;
 };
 env.wgpuWriteAtlasRegion = (handle, srcPtr, srcW, srcH, srcBytesPerRow, dstX, dstY, copyW, copyH) => {
    const tex = textures.get(handle);
    if (!tex) return;
    // For r8unorm, 1 byte per pixel; writeTexture requires bytesPerRow >= 256
    // OR == width if width*1 % 256 === 0 — for arbitrary widths we need to
    // re-pack into a 256-aligned staging buffer.
    const alignedBPR = Math.max(256, (srcBytesPerRow + 255) & ~255);
    if (alignedBPR === srcBytesPerRow) {
        const bytes = memU8().subarray(srcPtr + dstY * srcBytesPerRow + dstX,
                                       srcPtr + (dstY + copyH) * srcBytesPerRow);
        queue.writeTexture(
            { texture: tex, origin: { x: dstX, y: dstY } },
            bytes,
            { bytesPerRow: srcBytesPerRow, rowsPerImage: copyH },
            { width: copyW, height: copyH }
        );
    } else {
        // Repack copyW × copyH starting at (dstX, dstY) in the source.
        const staging = new Uint8Array(alignedBPR * copyH);
        const src = memU8();
        for (let y = 0; y < copyH; y++) {
            const srcRow = (dstY + y) * srcBytesPerRow + dstX;
            staging.set(src.subarray(srcPtr + srcRow, srcPtr + srcRow + copyW),
                        y * alignedBPR);
        }
        queue.writeTexture(
            { texture: tex, origin: { x: dstX, y: dstY } },
            staging,
            { bytesPerRow: alignedBPR, rowsPerImage: copyH },
            { width: copyW, height: copyH }
        );
    }
 };
 env.wgpuDestroyTexture = (handle) => {
    const tex = textures.get(handle);
    if (tex) { tex.destroy(); textures.delete(handle); textureViews.delete(handle); }
 };
 env.wgpuCreateLinearClampSampler = () => {
    const handle = newHandle();
    samplers.set(handle, device.createSampler({
        magFilter: "linear", minFilter: "linear",
        addressModeU: "clamp-to-edge", addressModeV: "clamp-to-edge",
    }));
    return handle;
 };
 // ─── per-frame ──────────────────────────────────────────────────────────
 env.wgpuFrameBegin = () => {
    state.frameBeginCount = (state.frameBeginCount || 0) + 1;
    if (state.gpuLost) return;
    ensureSized();
    state.encoder = device.createCommandEncoder();
    state.outIsPing = true;  // reset so each frame starts on the same target
    state.headerRingOffset = 0;
    // DON'T clearBuffer the header ring here. queue.writeBuffer ops from
    // writeHeader() are enqueued BEFORE this command buffer's submit,
    // so an encoded clearBuffer would wipe them — the dispatches would
    // then read all-zero headers and uiResolvePixel would reject every
    // pixel (surfaceW=0).
    clearStorageTexture(state.encoder, state.outIsPing ? state.pingTex : state.pongTex,
                        state.width, state.height);
    state.pass = state.encoder.beginComputePass();
 };
 let zeroBuffer = null;
 let zeroBufferSize = 0;
 function clearStorageTexture(encoder, tex, w, h) {
    const bpr = (w * 4 + 255) & ~255;
    const need = bpr * h;
    if (!zeroBuffer || zeroBufferSize < need) {
        if (zeroBuffer) zeroBuffer.destroy();
        zeroBuffer = device.createBuffer({ size: need, usage: GPUBufferUsage.COPY_SRC, mappedAtCreation: true });
        new Uint8Array(zeroBuffer.getMappedRange()).fill(0);
        zeroBuffer.unmap();
        zeroBufferSize = need;
    }
    encoder.copyBufferToTexture(
        { buffer: zeroBuffer, bytesPerRow: bpr, rowsPerImage: h },
        { texture: tex },
        { width: w, height: h, depthOrArrayLayers: 1 }
    );
 }
 env.wgpuFrameEnd = () => {
    state.frameEndCount = (state.frameEndCount || 0) + 1;
    if (state.gpuLost || !state.encoder) return;
    state.pass.end();
    state.pass = null;
    // Blit last-written ping-pong texture → canvas. After N dispatches,
    // state.outIsPing points at the NEXT write target, so the latest
    // content lives in the OPPOSITE texture.
    const finalTex = state.outIsPing ? state.pongTex : state.pingTex;
    const canvasTex = ctx.getCurrentTexture();
    state.encoder.copyTextureToTexture(
        { texture: finalTex },
        { texture: canvasTex },
        { width: state.width, height: state.height, depthOrArrayLayers: 1 }
    );
    queue.submit([state.encoder.finish()]);
    state.encoder = null;
 };
 // Write a 48-byte UIDispatchHeader into the ring buffer at the current
 // offset (which is incremented and 256-aligned). Returns the dynamic
 // offset to pass to setBindGroup.
 function writeHeader(headerPtr) {
    const offset = state.headerRingOffset;
    if (offset + HEADER_ALIGN > state.headerRingSize) {
        // Ring is small enough that overrun in one frame means too many
        // dispatches. Soft-wrap; correctness already requires the ring
        // be large enough.
        state.headerRingOffset = 0;
    }
    queue.writeBuffer(state.headerRing, state.headerRingOffset,
                      memU8().buffer, headerPtr, 48);
    state.headerRingOffset += HEADER_ALIGN;
    return offset;
 }
 function dispatchStandard(pipe, hdrBindGroup, headerPtr, gx, gy, itemsHandle, group3) {
    if (!state.pass) return;
    const off = writeHeader(headerPtr);
    state.pass.setPipeline(pipe.pipeline);
    state.pass.setBindGroup(0, hdrBindGroup, [off]);
    state.pass.setBindGroup(1, getGroup1BG(pipe.bgl1));
    state.pass.setBindGroup(2, getGroup2BG(pipe, itemsHandle));
    if (group3) state.pass.setBindGroup(3, group3);
    state.pass.dispatchWorkgroups(gx, gy, 1);
    // Flip ping-pong: the texture we just wrote becomes next dispatch's prev.
    state.outIsPing = !state.outIsPing;
 }
 env.wgpuDispatchQuads = (itemsHandle, headerPtr, gx, gy) => {
    state.dispatchQuadsCount = (state.dispatchQuadsCount || 0) + 1;
    dispatchStandard(pipeQuads, hdrBG.quads, headerPtr, gx, gy, itemsHandle, null);
 };
 env.wgpuDispatchCircles = (itemsHandle, headerPtr, gx, gy) => {
    dispatchStandard(pipeCircles, hdrBG.circles, headerPtr, gx, gy, itemsHandle, null);
 };
 env.wgpuDispatchImages = (itemsHandle, headerPtr, gx, gy, texHandle, sampHandle) => {
    const g3 = getGroup3BG(pipeImages, texHandle, sampHandle);
    dispatchStandard(pipeImages, hdrBG.images, headerPtr, gx, gy, itemsHandle, g3);
 };
 env.wgpuDispatchText = (itemsHandle, headerPtr, gx, gy, atlasHandle, sampHandle) => {
    const g3 = getGroup3BG(pipeText, atlasHandle, sampHandle);
    dispatchStandard(pipeText, hdrBG.text, headerPtr, gx, gy, itemsHandle, g3);
 };
 // ─── custom user-authored shaders ─────────────────────────────────────
 //
 // Bind-group contract (mirrors :WebGPUComputeShader.cppm):
 //   group 0 binding 0 — uniform UIDispatchHeader (dynamic offset, 48b)
 //   group 1 binding 0 — texture_storage_2d<rgba8unorm, write> out
 //   group 1 binding 1 — texture_2d<f32>                     prev
 //   group 2+          — user-declared (UICustomBinding entries)
 //
 // Each UICustomBinding entry on the wasm side is 8 bytes:
 //   u8 group, u8 binding, u8 kind, u8 pad, u32 pushOffset
 // kind: 0 = read-only-storage SSBO, 1 = sampled tex 2d, 2 = filtering sampler.
 const customPipelines = new Map();   // handle → { pipeline, bgls, hdrBG, byGroup }
 env.wgpuLoadCustomShader = (wgslPtr, wgslLen, bindingsPtr, bindingsCount) => {
    const wgsl = new TextDecoder().decode(memU8().subarray(wgslPtr, wgslPtr + wgslLen));
    const bindings = [];
    const dv = new DataView(memU8().buffer, bindingsPtr, bindingsCount * 8);
    for (let i = 0; i < bindingsCount; i++) {
        bindings.push({
            group: dv.getUint8(i*8 + 0),
            binding: dv.getUint8(i*8 + 1),
            kind: dv.getUint8(i*8 + 2),
            pushOffset: dv.getUint32(i*8 + 4, true),
        });
    }
    // Group bindings by @group(N) for layout creation.
    const byGroup = new Map();
    for (const b of bindings) {
        if (b.group < 2) {
            console.error(`[crafter-wgpu] custom shader: @group(${b.group}) reserved; use groups >= 2`);
            return 0;
        }
        if (!byGroup.has(b.group)) byGroup.set(b.group, []);
        byGroup.get(b.group).push(b);
    }
    // Group 0 = header uniform, Group 1 = ping-pong out+prev — always injected.
    const bgls = [
        device.createBindGroupLayout({ entries: [
            { binding: 0, visibility: GPUShaderStage.COMPUTE,
              buffer: { type: "uniform", hasDynamicOffset: true, minBindingSize: 48 } },
        ]}),
        device.createBindGroupLayout({ entries: [
            { binding: 0, visibility: GPUShaderStage.COMPUTE,
              storageTexture: { format: "rgba8unorm", access: "write-only", viewDimension: "2d" } },
            { binding: 1, visibility: GPUShaderStage.COMPUTE,
              texture: { sampleType: "float", viewDimension: "2d" } },
        ]}),
    ];
    // Sorted custom groups. Pad any gaps with empty bgls (WebGPU pipeline
    // layouts require a contiguous array of GPUBindGroupLayout per group
    // index up to the highest used).
    const sortedGroups = [...byGroup.keys()].sort((a, b) => a - b);
    const highest = sortedGroups.length ? sortedGroups[sortedGroups.length - 1] : 1;
    for (let g = 2; g <= highest; g++) {
        if (byGroup.has(g)) {
            const entries = byGroup.get(g).map(b => {
                const e = { binding: b.binding, visibility: GPUShaderStage.COMPUTE };
                if      (b.kind === 0) e.buffer = { type: "read-only-storage" };
                else if (b.kind === 1) e.texture = { sampleType: "float", viewDimension: "2d" };
                else if (b.kind === 2) e.sampler = { type: "filtering" };
                return e;
            });
            bgls.push(device.createBindGroupLayout({ entries }));
        } else {
            bgls.push(device.createBindGroupLayout({ entries: [] }));
        }
    }
    let pipeline;
    try {
        const mod = device.createShaderModule({ code: wgsl });
        const layout = device.createPipelineLayout({ bindGroupLayouts: bgls });
        pipeline = device.createComputePipeline({ layout, compute: { module: mod, entryPoint: "main" } });
    } catch (e) {
        console.error("[crafter-wgpu] custom shader compile failed:", e);
        return 0;
    }
    const hdrBG = device.createBindGroup({
        layout: bgls[0],
        entries: [{ binding: 0, resource: { buffer: state.headerRing, offset: 0, size: 48 } }],
    });
    const handle = newHandle();
    customPipelines.set(handle, { pipeline, bgls, hdrBG, byGroup, sortedGroups });
    return handle;
 };
 env.wgpuDispatchCustom = (pipelineHandle, pushPtr, pushBytes, handlesPtr, handlesCount,
                          gx, gy, gz) => {
    state.dispatchCustomCount = (state.dispatchCustomCount || 0) + 1;
    if (!state.pass) return;
    const pipe = customPipelines.get(pipelineHandle);
    if (!pipe) {
        console.error("[crafter-wgpu] wgpuDispatchCustom: unknown pipeline", pipelineHandle);
        return;
    }
    // Write header (first 48 bytes of push).
    const off = writeHeader(pushPtr);
    state.pass.setPipeline(pipe.pipeline);
    state.pass.setBindGroup(0, pipe.hdrBG, [off]);
    state.pass.setBindGroup(1, getGroup1BG(pipe.bgls[1]));
    // Walk bindings in declaration order and assemble bind groups.
    // handles[] from wasm is in the SAME order as customBindings, so we
    // pick up indices by walking byGroup in the same sorted order.
    const handles = new Uint32Array(memU8().buffer, handlesPtr, handlesCount);
    let handleIdx = 0;
    for (const g of pipe.sortedGroups) {
        const entries = pipe.byGroup.get(g).map(b => {
            const h = handles[handleIdx++];
            let resource;
            if      (b.kind === 0) resource = { buffer: buffers.get(h) };
            else if (b.kind === 1) resource = textureViews.get(h);
            else if (b.kind === 2) resource = samplers.get(h);
            return { binding: b.binding, resource };
        });
        const bg = device.createBindGroup({ layout: pipe.bgls[g], entries });
        state.pass.setBindGroup(g, bg);
    }
    state.pass.dispatchWorkgroups(gx, gy, gz);
    state.outIsPing = !state.outIsPing;
 };
 // Debug accessor for browser-console diagnostics.
 window.crafter_wgpu_state = state;
 window.crafter_wgpu_device = device;
 window.crafter_wgpu_canvasCtx = ctx;
 window.crafter_wgpu_debug = () => ({
    width: state.width, height: state.height,
    outIsPing: state.outIsPing,
    encoderActive: !!state.encoder,
    passActive: !!state.pass,
    bgCacheSize: state.bindGroupCache.size,
    bufferHandles: buffers.size,
    textureHandles: textures.size,
    samplerHandles: samplers.size,
    headerRingOffset: state.headerRingOffset,
    frameBeginCount: state.frameBeginCount || 0,
    frameEndCount: state.frameEndCount || 0,
    dispatchQuadsCount: state.dispatchQuadsCount || 0,
    writeBufferCount: state.writeBufferCount || 0,
    lastWriteHandle: state.lastWriteHandle,
    lastWriteSize: state.lastWriteSize,
 });
 window.crafter_wgpu_bufferKeys = () => [...buffers.keys()];
 // Read back the first QuadItem from a registered buffer to verify the
 // GPU sees what the CPU wrote.
 window.crafter_wgpu_readBuffer = async (handle, byteSize = 64) => {
    const buf = buffers.get(handle);
    if (!buf) return "no buffer for handle " + handle;
    const read = device.createBuffer({ size: 256, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST });
    const enc = device.createCommandEncoder();
    enc.copyBufferToBuffer(buf, 0, read, 0, byteSize);
    device.queue.submit([enc.finish()]);
    await read.mapAsync(GPUMapMode.READ);
    const data = new Float32Array(read.getMappedRange().slice(0, byteSize));
    read.unmap();
    return [...data];
 };
 // Surface size getters (the wasm side may query these on Resize events).
 env.wgpuSurfaceWidth  = () => state.width || canvas.width;
 env.wgpuSurfaceHeight = () => state.height || canvas.height;
 // One-shot init: forces ping-pong allocation at current canvas size so
 // any Buffer/Texture creation before the first frame works against a
 // concrete size. Called by Crafter::Device::Initialize on the wasm side.
 env.wgpuInit = () => {
    const { w, h } = syncCanvasSize();
    recreatePingPong(w, h);
 };
 // Resize listener — wires up to the same `resize` event dom-env.js
 // listens to. We trigger sizing on next frame begin; no work here.
 window.addEventListener("resize", () => { /* ensureSized in wgpuFrameBegin */ });
 console.log("[crafter-wgpu] init complete; env handlers wired");
 } catch (e) {
    // Capture any throw so the stub error messages name the real cause
    // instead of "(no error captured)". Re-throw so runtime.js's catch
    // also logs it.
    initError = e instanceof Error ? e : new Error(String(e));
    console.error("[crafter-wgpu] init failed:", initError);
    throw initError;
 }
 })(); // end window.crafter_webbuild_env_ready
--- a/examples/CustomShader/inverse-circle.comp.wgsl
+++ b/examples/CustomShader/inverse-circle.comp.wgsl
@ -1,58 +0,0 @@
 // DOM-mode port of inverse-circle.comp.glsl. Inverts RGB inside each
 // user-supplied circle; passes through every other pixel so the
 // ping-pong carries the prior dispatch's scene forward.
 //
 // Bind-group contract (mirrors :WebGPUComputeShader.cppm):
 //   group 0 binding 0 — uniform UIDispatchHeader (auto-injected)
 //   group 1 binding 0 — texture_storage_2d<rgba8unorm, write> out (auto)
 //   group 1 binding 1 — texture_2d<f32>                      prev (auto)
 //   group 2 binding 0 — items SSBO declared by the C++ side via
 //                       UICustomBinding { group=2, binding=0,
 //                       kind=Buffer, pushOffset=offsetof(hdr.itemBuffer) }
 struct UIDispatchHeader {
    outImage:   u32,
    itemBuffer: u32,
    surfaceW:   u32,
    surfaceH:   u32,
    clipX:      f32,
    clipY:      f32,
    clipW:      f32,
    clipH:      f32,
    itemCount:  u32,
    frameIdx:   u32,
    flags:      u32,
    _pad:       u32,
 };
@group(0) @binding(0) var<uniform> hdr : UIDispatchHeader;
@group(1) @binding(0) var outTex  : texture_storage_2d<rgba8unorm, write>;
@group(1) @binding(1) var prevTex : texture_2d<f32>;
 struct InverseCircleItem {
    centerRadius: vec4<f32>,
 };
@group(2) @binding(0) var<storage, read> items : array<InverseCircleItem>;
@compute @workgroup_size(8, 8, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
    if (gid.x >= hdr.surfaceW || gid.y >= hdr.surfaceH) { return; }
    let coord = vec2<i32>(i32(gid.x), i32(gid.y));
    let sp = vec2<f32>(f32(gid.x) + 0.5, f32(gid.y) + 0.5);
    // Max-coverage over all items, so overlapping circles don't
    // double-invert back to the original.
    var coverage: f32 = 0.0;
    for (var i: u32 = 0u; i < hdr.itemCount; i = i + 1u) {
        let it = items[i];
        let c = it.centerRadius.xy;
        let r = it.centerRadius.z;
        let d = length(sp - c) - r;
        coverage = max(coverage, clamp(0.5 - d, 0.0, 1.0));
    }
    var dst = textureLoad(prevTex, coord, 0);
    if (coverage > 0.0) {
        dst = vec4<f32>(mix(dst.rgb, vec3<f32>(1.0) - dst.rgb, coverage), dst.a);
    }
    textureStore(outTex, coord, dst);
 }
--- a/examples/CustomShader/main.cpp
+++ b/examples/CustomShader/main.cpp
@ -2,17 +2,8 @@
 // standard ones. The custom shader inverts RGB in the area covered by a
 // list of circles. The mouse-tracking circle moves; two static ones sit
 // on a striped background drawn with the standard DrawQuads shader.
 //
 // Works on both Vulkan (native) and WebGPU (DOM). The shader source is
 // in two files — inverse-circle.comp.glsl (native, SPIR-V) and
 // inverse-circle.comp.wgsl (DOM). The C++ surface differs only at the
 // shader-load / buffer-flag sites; the per-frame UI building logic is
 // identical.
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
 #endif
 #include <cstddef>  // offsetof — a macro, so not visible via `import std;`
 import Crafter.Graphics;
 import Crafter.Event;
@ -20,22 +11,18 @@ import std;
 using namespace Crafter;
 // Application-side item POD. Matches `struct InverseCircleItem { vec4
-// centerRadius; }` in inverse-circle.comp.{glsl,wgsl} byte-for-byte.
+// centerRadius; }` in inverse-circle.comp.glsl byte-for-byte.
 struct InverseCircleItem {
    float cx, cy, radius, _pad;
 };
 int main() {
    Device::Initialize();
 #ifdef CRAFTER_GRAPHICS_WINDOW_DOM
    static Window window(1280, 720, "Custom Shader");
 #else
    Window window(1280, 720, "Custom Shader");
 #endif
-    auto init = window.StartInit();
+    VkCommandBuffer init = window.StartInit();
-    GraphicsDescriptorHeap heap;
+    DescriptorHeapVulkan heap;
    heap.Initialize(/*images*/ 8, /*buffers*/ 8, /*samplers*/ 4);
    window.descriptorHeap = &heap;
@ -43,46 +30,26 @@ int main() {
    ui.Initialize(window, heap, init);
    window.passes.push_back(&ui);
-    // Load the user-authored shader. On native it's an offline-compiled
+    // Load the user-authored shader. Same wrapper as the four shipped with
-    // SPIR-V from .comp.glsl; on DOM it's WGSL source compiled at
+    // the library — there is no privileged path.
-    // startup by the device. The DOM Load takes an extra `bindings` arg
+    ComputeShader inverseCircle;
    // declaring resources the renderer should bind at dispatch time —
    // here just one entry: the items SSBO at group 2, with its heap slot
    // read from `hdr.itemBuffer` in the push data.
    GraphicsComputeShader inverseCircle;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
    inverseCircle.Load("inverse-circle.comp.spv");
 #else
    UICustomBinding invBindings[] = {
        { .group   = 2,
          .binding = 0,
          .kind    = UICustomBindingKind::Buffer,
          ._pad    = 0,
          .pushOffset = static_cast<std::uint32_t>(offsetof(Crafter::UIDispatchHeader, itemBuffer)) },
    };
    inverseCircle.Load(std::filesystem::path("inverse-circle.comp.wgsl"), invBindings);
 #endif
    // User-owned buffers.
-    GraphicsBuffer<QuadItem,          true> quadsBuf;
+    VulkanBuffer<QuadItem,           true> quadsBuf;
-    GraphicsBuffer<InverseCircleItem, true> invBuf;
+    VulkanBuffer<InverseCircleItem,  true> invBuf;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
    quadsBuf.Create(
        VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT,
        VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, 64);
    invBuf.Create(
        VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT,
        VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, 16);
 #else
    quadsBuf.Create(64);
    invBuf.Create(16);
 #endif
    auto quadsSlot = ui.RegisterBuffer(quadsBuf);
    auto invSlot   = ui.RegisterBuffer(invBuf);
    EventListener<UIBuildArgs> buildSub(&ui.onBuild, [&](UIBuildArgs a) {
-        auto cmd = a.cmd;
+        VkCommandBuffer cmd = a.cmd;
        Rect canvas = Rect::FromWindow(window);
@ -115,24 +82,15 @@ int main() {
        // Standard dispatch first — paints the stripes.
        if (qc > 0) {
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
            quadsBuf.FlushDevice(cmd, VK_ACCESS_SHADER_READ_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT);
 #else
            quadsBuf.FlushDevice();
 #endif
            ui.DispatchQuads(cmd, quadsSlot, qc);
        }
        // Custom dispatch second — reads the stripes, inverts under
        // circles, writes back. The library inserts the inter-dispatch
-        // SHADER_WRITE → SHADER_READ|WRITE barrier automatically on
+        // SHADER_WRITE → SHADER_READ|WRITE barrier automatically.
        // Vulkan; WebGPU's compute-pass tracking does it for free.
        if (ic > 0) {
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
            invBuf.FlushDevice(cmd, VK_ACCESS_SHADER_READ_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT);
 #else
            invBuf.FlushDevice();
 #endif
            struct PC { UIDispatchHeader hdr; } pc { ui.FillHeader(invSlot, ic) };
            std::uint32_t gx = (window.width  + 7) / 8;
            std::uint32_t gy = (window.height + 7) / 8;
--- a/examples/CustomShader/project.cpp
+++ b/examples/CustomShader/project.cpp
@ -4,14 +4,6 @@ namespace fs = std::filesystem;
 using namespace Crafter;
 extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> args) {
    bool isWasm = false;
    for (std::string_view a : args) {
        if (a.starts_with("--target=") && a.find("wasm") != std::string_view::npos) {
            isWasm = true;
            break;
        }
    }
    Configuration* graphics = LocalProject({
        .projectFile = "../../project.cpp",
        .args = std::vector<std::string>(args.begin(), args.end()),
@ -21,14 +13,6 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
    cfg.path = "./";
    cfg.name = "CustomShader";
    cfg.outputName = "CustomShader";
    cfg.type = ConfigurationType::Executable;
    if (isWasm) {
        cfg.target = "wasm32-wasip1";
        cfg.defines.push_back({"CRAFTER_GRAPHICS_WINDOW_DOM", ""});
        // Match the -msimd128 that Crafter.Math/Crafter.Graphics's wasm
        // PCMs were compiled with.
        cfg.compileFlags.push_back("-msimd128");
    }
    ApplyStandardArgs(cfg, args);
    cfg.dependencies = { graphics };
@ -36,14 +20,6 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
    std::array<fs::path, 1> impls = { "main" };
    cfg.GetInterfacesAndImplementations(ifaces, impls);
-    if (isWasm) {
+    cfg.shaders.emplace_back(fs::path("inverse-circle.comp.glsl"), std::string("main"), ShaderType::Compute);
        // WGSL source is shipped as a static file and loaded via the
        // WASI VFS at runtime through WebGPUComputeShader::Load(path).
        cfg.files.emplace_back(fs::path("inverse-circle.comp.wgsl"));
        EnableWasiBrowserRuntime(cfg);
    } else {
        cfg.shaders.emplace_back(fs::path("inverse-circle.comp.glsl"),
                                 std::string("main"), ShaderType::Compute);
    }
    return cfg;
 }
--- a/examples/HelloDom/serve.sh
+++ b/examples/HelloDom/serve.sh
@ -0,0 +1,2 @@
 #!/usr/bin/env sh
 caddy file-server --listen :8080 --root bin/HelloDom-wasm32-wasip1-native-native-df37fe0fe124fe57
--- a/examples/HelloUI/main.cpp
+++ b/examples/HelloUI/main.cpp
@ -4,9 +4,7 @@
 // (Tier 2 standard shader, dispatched directly). Hit-testing for the button
 // and slider is the user's responsibility — see the onMouseMove listener.
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
 #endif
 import Crafter.Graphics;
 import Crafter.Event;
@ -15,15 +13,11 @@ using namespace Crafter;
 int main() {
    Device::Initialize();
    #ifdef CRAFTER_GRAPHICS_WINDOW_DOM
    static Window window(1280, 720, "Hello UI");
    #else
    Window window(1280, 720, "Hello UI");
    #endif
-    auto init = window.StartInit();
+    VkCommandBuffer init = window.StartInit();
-    GraphicsDescriptorHeap heap;
+    DescriptorHeapVulkan heap;
    heap.Initialize(/*images*/ 16, /*buffers*/ 16, /*samplers*/ 4);
    window.descriptorHeap = &heap;
@ -39,10 +33,9 @@ int main() {
    // User-owned per-shader buffers. Mapped, written each frame, dispatched
    // by the user. Capacity is up to the user; resize means re-Register.
-    GraphicsBuffer<QuadItem,   true> quadsBuf;
+    VulkanBuffer<QuadItem,   true> quadsBuf;
-    GraphicsBuffer<CircleItem, true> circlesBuf;
+    VulkanBuffer<CircleItem, true> circlesBuf;
-    GraphicsBuffer<GlyphItem,  true> glyphsBuf;
+    VulkanBuffer<GlyphItem,  true> glyphsBuf;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
    quadsBuf.Create(
        VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT,
        VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, 256);
@ -52,11 +45,6 @@ int main() {
    glyphsBuf.Create(
        VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT,
        VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, 4096);
 #else
    quadsBuf.Create(256);
    circlesBuf.Create(64);
    glyphsBuf.Create(4096);
 #endif
    auto quadsSlot   = ui.RegisterBuffer(quadsBuf);
    auto circlesSlot = ui.RegisterBuffer(circlesBuf);
@ -114,7 +102,7 @@ int main() {
    };
    EventListener<UIBuildArgs> buildSub(&ui.onBuild, [&](UIBuildArgs a) {
-        auto cmd = a.cmd;
+        VkCommandBuffer cmd = a.cmd;
        // Update demo progress.
        progress = std::fmod(progress + 0.005f, 1.0f);
@ -172,27 +160,15 @@ int main() {
        // Flush + dispatch. The library inserts the inter-dispatch barriers.
        if (qc > 0) {
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
            quadsBuf.FlushDevice(cmd, VK_ACCESS_SHADER_READ_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT);
 #else
            quadsBuf.FlushDevice();
 #endif
            ui.DispatchQuads(cmd, quadsSlot, qc);
        }
        if (cc > 0) {
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
            circlesBuf.FlushDevice(cmd, VK_ACCESS_SHADER_READ_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT);
 #else
            circlesBuf.FlushDevice();
 #endif
            ui.DispatchCircles(cmd, circlesSlot, cc);
        }
        if (gc > 0) {
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
            glyphsBuf.FlushDevice(cmd, VK_ACCESS_SHADER_READ_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT);
 #else
            glyphsBuf.FlushDevice();
 #endif
            ui.DispatchText(cmd, glyphsSlot, gc);
        }
    });
--- a/examples/HelloUI/project.cpp
+++ b/examples/HelloUI/project.cpp
@ -4,14 +4,6 @@ namespace fs = std::filesystem;
 using namespace Crafter;
 extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> args) {
    bool isWasm = false;
    for (std::string_view a : args) {
        if (a.starts_with("--target=") && a.find("wasm") != std::string_view::npos) {
            isWasm = true;
            break;
        }
    }
    Configuration* graphics = LocalProject({
        .projectFile = "../../project.cpp",
        .args = std::vector<std::string>(args.begin(), args.end()),
@ -21,15 +13,6 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
    cfg.path = "./";
    cfg.name = "HelloUI";
    cfg.outputName = "HelloUI";
    cfg.type = ConfigurationType::Executable;
    if (isWasm) {
        cfg.target = "wasm32-wasip1";
        cfg.defines.push_back({"CRAFTER_GRAPHICS_WINDOW_DOM", ""});
        // Match the -msimd128 that Crafter.Math/Crafter.Graphics's wasm
        // PCMs were compiled with, otherwise PCM imports trip a config-
        // mismatch error in recent clangs.
        cfg.compileFlags.push_back("-msimd128");
    }
    ApplyStandardArgs(cfg, args);
    cfg.dependencies = { graphics };
@ -38,9 +21,5 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
    cfg.GetInterfacesAndImplementations(ifaces, impls);
    cfg.files.push_back("font.ttf");
    if (isWasm) {
        EnableWasiBrowserRuntime(cfg);
    }
    return cfg;
 }
--- a/implementations/Crafter.Graphics-FontAtlas.cpp
+++ b/implementations/Crafter.Graphics-FontAtlas.cpp
@ -2,53 +2,47 @@
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 This library is free software; you can redistribute it and/or
 modify it under the terms of the GNU Lesser General Public
 License version 3.0 as published by the Free Software Foundation;
 This library is distributed in the hope that it will be useful,
 but WITHOUT ANY WARRANTY; without even the implied warranty of
 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
 Lesser General Public License for more details.
 You should have received a copy of the GNU Lesser General Public
 License along with this library; if not, write to the Free Software
 Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301 USA
 */
 module;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
 #endif
 #include "../lib/stb_truetype.h"
 module Crafter.Graphics:FontAtlas_impl;
 import :FontAtlas;
 import :Font;
 import :GraphicsTypes;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import :ImageVulkan;
 import :Device;
 #else
 import :WebGPU;
 #endif
 import std;
 using namespace Crafter;
-std::uint8_t* FontAtlas::PixelPtr() noexcept {
+void FontAtlas::Initialize(VkCommandBuffer cmd) {
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
    return image.buffer.value;
 #else
    return staging.data();
 #endif
 }
 void FontAtlas::Initialize(GraphicsCommandBuffer cmd) {
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
    image.Create(
        kAtlasSize, kAtlasSize, /*mipLevels*/ 1, cmd,
        VK_FORMAT_R8_UNORM,
        VK_IMAGE_USAGE_TRANSFER_DST_BIT | VK_IMAGE_USAGE_SAMPLED_BIT,
        VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL
    );
    // Staging buffer is mapped; clear it so empty atlas regions sample as
    // distance < onedge (i.e. fully outside any glyph).
    std::memset(image.buffer.value, 0, kAtlasSize * kAtlasSize);
    dirty = true;
 #else
    (void)cmd;
    staging.assign(kAtlasSize * kAtlasSize, 0);
    textureHandle = WebGPU::wgpuCreateAtlasTexture(kAtlasSize, kAtlasSize);
    dirty = true;
 #endif
 }
 bool FontAtlas::ShelfPlace(int w, int h, int& outX, int& outY) {
    // Try existing shelves first — same height heuristic keeps fragmentation low.
    for (Shelf& s : shelves_) {
        if (h <= s.height && s.cursorX + w <= kAtlasSize) {
            outX = s.cursorX;
@ -57,6 +51,7 @@ bool FontAtlas::ShelfPlace(int w, int h, int& outX, int& outY) {
            return true;
        }
    }
    // New shelf below current ones.
    if (nextShelfY_ + h > kAtlasSize) return false;
    Shelf s{};
    s.y       = nextShelfY_;
@ -75,6 +70,7 @@ bool FontAtlas::Ensure(Font& font, std::uint32_t codepoint) {
    float fontScale = stbtt_ScaleForPixelHeight(&font.font, kBaseSize);
    // Advance is always present, even for empty glyphs (e.g. space).
    int advanceUnits = 0, lsb = 0;
    stbtt_GetCodepointHMetrics(&font.font, static_cast<int>(codepoint), &advanceUnits, &lsb);
@ -94,12 +90,12 @@ bool FontAtlas::Ensure(Font& font, std::uint32_t codepoint) {
        int px = 0, py = 0;
        if (!ShelfPlace(sw, sh, px, py)) {
            stbtt_FreeSDF(sdf, nullptr);
-            return false;
+            return false;       // V1: silently drop overflow; V2: grow atlas
        }
-        std::uint8_t* dst = PixelPtr();
+        // Blit row-by-row into the mapped staging buffer.
        for (int row = 0; row < sh; ++row) {
            std::memcpy(
-                dst + (py + row) * kAtlasSize + px,
+                image.buffer.value + (py + row) * kAtlasSize + px,
                sdf + row * sw,
                static_cast<std::size_t>(sw)
            );
@ -114,6 +110,8 @@ bool FontAtlas::Ensure(Font& font, std::uint32_t codepoint) {
        g.v1 = static_cast<float>(py + sh) / kAtlasSize;
        dirty = true;
    }
    // For empty glyphs (whitespace) we still cache the entry — the size-0
    // fields tell the emitter to skip the quad but advance the cursor.
    cache_.emplace(key, g);
    return true;
@ -124,18 +122,8 @@ const Glyph* FontAtlas::Lookup(Font& font, std::uint32_t codepoint) const {
    return it == cache_.end() ? nullptr : &it->second;
 }
-void FontAtlas::Update(GraphicsCommandBuffer cmd) {
+void FontAtlas::Update(VkCommandBuffer cmd) {
    if (!dirty) return;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
    image.Update(cmd, VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL);
 #else
    (void)cmd;
    // Full-atlas upload. Future: track dirty region.
    WebGPU::wgpuWriteAtlasRegion(
        textureHandle, staging.data(),
        kAtlasSize, kAtlasSize, kAtlasSize,
        0, 0, kAtlasSize, kAtlasSize
    );
 #endif
    dirty = false;
 }
--- a/implementations/Crafter.Graphics-UI-Shared.cpp
+++ b/implementations/Crafter.Graphics-UI-Shared.cpp
@ -1,86 +0,0 @@
 /*
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 // Backend-agnostic UIRenderer methods. FillHeader and ShapeText do not
 // touch any GPU API; they only read window dimensions and (for ShapeText)
 // the CPU-side font atlas. Split out so both Vulkan and WebGPU impls
 // share the same source.
 module Crafter.Graphics:UI_shared_impl;
 import :UI;
 import :Window;
 import :Font;
 import :FontAtlas;
 import :GraphicsTypes;
 import std;
 using namespace Crafter;
 UIDispatchHeader UIRenderer::FillHeader(std::uint32_t itemBufferSlot,
                                        std::uint32_t itemCount,
                                        std::array<float,4> clipRectPx,
                                        std::uint32_t flags) const noexcept {
    UIDispatchHeader h{};
    h.outImage      = outImageSlot_;
    h.itemBuffer    = itemBufferSlot;
    h.surfaceWidth  = window_->width;
    h.surfaceHeight = window_->height;
    h.clipX = clipRectPx[0];
    h.clipY = clipRectPx[1];
    h.clipW = clipRectPx[2];
    h.clipH = clipRectPx[3];
    h.itemCount = itemCount;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
    h.frameIdx  = window_->currentBuffer;
 #else
    h.frameIdx  = 0;
 #endif
    h.flags     = flags;
    h._pad      = 0;
    return h;
 }
 std::uint32_t UIRenderer::ShapeText(Font& font, float pxSize,
                                    float x, float baselineY,
                                    std::string_view utf8,
                                    std::array<float,4> color,
                                    GlyphItem* out, std::uint32_t outCapacity,
                                    float* outAdvance) {
    if (fontAtlas == nullptr) {
        std::println("UIRenderer::ShapeText: no FontAtlas (set fontAtlas before Initialize)");
        std::abort();
    }
    const float scale = pxSize / FontAtlas::kBaseSize;
    float cursor = x;
    std::uint32_t written = 0;
    std::size_t i = 0;
    while (i < utf8.size() && written < outCapacity) {
        std::uint32_t cp = DecodeUtf8(utf8, i);
        if (cp == 0) break;
        if (cp == '\n') { continue; }
        fontAtlas->Ensure(font, cp);
        const Glyph* g = fontAtlas->Lookup(font, cp);
        if (g == nullptr) continue;
        if (g->w > 0 && g->h > 0) {
            GlyphItem& gi = out[written++];
            gi.x  = cursor + g->xoff * scale;
            gi.y  = baselineY + g->yoff * scale;
            gi.w  = g->w * scale;
            gi.h  = g->h * scale;
            gi.u0 = g->u0; gi.v0 = g->v0;
            gi.u1 = g->u1; gi.v1 = g->v1;
            gi.r = color[0]; gi.g = color[1]; gi.b = color[2]; gi.a = color[3];
        }
        cursor += g->advance * scale;
    }
    if (outAdvance) *outAdvance = cursor - x;
    return written;
 }
--- a/implementations/Crafter.Graphics-UI-WebGPU.cpp
+++ b/implementations/Crafter.Graphics-UI-WebGPU.cpp
@ -1,180 +0,0 @@
 /*
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 // WebGPU UIRenderer implementation — DOM mode parallel to
 // Crafter.Graphics-UI.cpp. Compute pipelines and bind groups live JS-side
 // in additional/dom-webgpu.js; this file just translates UIRenderer's
 // public method calls into wgpu* import calls.
 module Crafter.Graphics:UI_webgpu_impl;
 import :UI;
 import :Window;
 import :Font;
 import :FontAtlas;
 import :WebGPU;
 import :WebGPUBuffer;
 import :DescriptorHeapWebGPU;
 import :GraphicsTypes;
 import std;
 using namespace Crafter;
 void UIRenderer::Initialize(Window& window, GraphicsDescriptorHeap& heap, GraphicsCommandBuffer initCmd,
                            std::filesystem::path /*quadsSpv*/,
                            std::filesystem::path /*circlesSpv*/,
                            std::filesystem::path /*imagesSpv*/,
                            std::filesystem::path /*textSpv*/) {
    (void)initCmd;
    window_ = &window;
    heap_   = &heap;
    // The JS bridge owns the compute pipelines (4 of them, precompiled at
    // page load). The C++ side has nothing to load. We do reserve one
    // image slot for the swapchain output (kept symmetrical with Vulkan)
    // and register the font atlas if one was set.
    auto outRange = heap_->AllocateImageSlots(1);
    outImageSlot_ = ImageSlot{heap_, outRange.firstElement};
    // No JS texture handle stored at this slot — the canvas's
    // getCurrentTexture() is per-frame and the JS bridge resolves it
    // internally. The slot value is only meaningful through
    // UIDispatchHeader.outImage, which the WGSL shaders ignore.
    if (fontAtlas != nullptr) {
        auto atlasImg = heap_->AllocateImageSlots(1);
        fontAtlasImageSlot_ = ImageSlot{heap_, atlasImg.firstElement};
        heap_->imageTable[atlasImg.firstElement] = fontAtlas->textureHandle;
        fontAtlasSamplerSlot_ = RegisterLinearClampSampler();
    }
    resizeSub_.SetEvent(&window.onResize, [this]() {
        // Storage textures + bind groups are recreated JS-side at next
        // wgpuFrameBegin (via the ensureSized path). Nothing to do here.
    });
 }
 void UIRenderer::Record(GraphicsCommandBuffer cmd, std::uint32_t frameIdx, Window& window) {
    if (fontAtlas != nullptr && fontAtlas->dirty) {
        fontAtlas->Update(cmd);
    }
    onBuild.Invoke({cmd, frameIdx});
    if (fontAtlas != nullptr && fontAtlas->dirty) {
        fontAtlas->Update(cmd);
    }
    (void)window;
 }
 namespace {
    inline std::uint32_t TilesFor(std::uint32_t dim) { return (dim + 7u) / 8u; }
 }
 void UIRenderer::DispatchQuads(GraphicsCommandBuffer /*cmd*/, std::uint32_t bufferSlot,
                               std::uint32_t itemCount,
                               std::array<float,4> clipRectPx) {
    if (itemCount == 0) return;
    UIDispatchHeader hdr = FillHeader(bufferSlot, itemCount, clipRectPx);
    auto handle = heap_->bufferTable[bufferSlot];
    WebGPU::wgpuDispatchQuads(handle, &hdr,
        static_cast<std::int32_t>(TilesFor(window_->width)),
        static_cast<std::int32_t>(TilesFor(window_->height)));
 }
 void UIRenderer::DispatchCircles(GraphicsCommandBuffer /*cmd*/, std::uint32_t bufferSlot,
                                 std::uint32_t itemCount,
                                 std::array<float,4> clipRectPx) {
    if (itemCount == 0) return;
    UIDispatchHeader hdr = FillHeader(bufferSlot, itemCount, clipRectPx);
    auto handle = heap_->bufferTable[bufferSlot];
    WebGPU::wgpuDispatchCircles(handle, &hdr,
        static_cast<std::int32_t>(TilesFor(window_->width)),
        static_cast<std::int32_t>(TilesFor(window_->height)));
 }
 void UIRenderer::DispatchImages(GraphicsCommandBuffer /*cmd*/, std::uint32_t bufferSlot,
                                std::uint32_t itemCount,
                                std::array<float,4> clipRectPx) {
    if (itemCount == 0) return;
    UIDispatchHeader hdr = FillHeader(bufferSlot, itemCount, clipRectPx);
    auto handle = heap_->bufferTable[bufferSlot];
    // For DispatchImages, the WGSL expects a texture + sampler in group 3.
    // The library v1 doesn't expose user-image registration on DOM (out of
    // scope per plan). If the user calls DispatchImages without a registered
    // image, fall back to using the font atlas binding — the user's items
    // should reference texSlot/sampSlot but on DOM those are ignored. For
    // now, route through the font atlas texture if available; otherwise
    // skip the dispatch.
    if (fontAtlasImageSlot_) {
        auto texHandle  = heap_->imageTable[fontAtlasImageSlot_];
        auto sampHandle = heap_->samplerTable[fontAtlasSamplerSlot_];
        WebGPU::wgpuDispatchImages(handle, &hdr,
            static_cast<std::int32_t>(TilesFor(window_->width)),
            static_cast<std::int32_t>(TilesFor(window_->height)),
            texHandle, sampHandle);
    }
 }
 void UIRenderer::DispatchText(GraphicsCommandBuffer /*cmd*/, std::uint32_t bufferSlot,
                              std::uint32_t itemCount,
                              std::array<float,4> clipRectPx) {
    if (itemCount == 0) return;
    if (!fontAtlasImageSlot_) {
        std::println("UIRenderer::DispatchText: no FontAtlas registered (set fontAtlas before Initialize)");
        std::abort();
    }
    UIDispatchHeader hdr = FillHeader(bufferSlot, itemCount, clipRectPx);
    auto bufHandle  = heap_->bufferTable[bufferSlot];
    auto texHandle  = heap_->imageTable[fontAtlasImageSlot_];
    auto sampHandle = heap_->samplerTable[fontAtlasSamplerSlot_];
    WebGPU::wgpuDispatchText(bufHandle, &hdr,
        static_cast<std::int32_t>(TilesFor(window_->width)),
        static_cast<std::int32_t>(TilesFor(window_->height)),
        texHandle, sampHandle);
 }
 SamplerSlot UIRenderer::RegisterLinearClampSampler() {
    auto range = heap_->AllocateSamplerSlots(1);
    heap_->samplerTable[range.firstElement] = WebGPU::wgpuCreateLinearClampSampler();
    return SamplerSlot{heap_, range.firstElement};
 }
 void UIRenderer::Dispatch(GraphicsCommandBuffer /*cmd*/, const GraphicsComputeShader& shader,
                          const void* push, std::uint32_t pushBytes,
                          std::uint32_t gx, std::uint32_t gy, std::uint32_t gz) {
    // For each user-declared binding, read the slot uint32 out of push
    // data at the recorded offset, look up the GPU handle in the heap,
    // and assemble a list of handles in the same order the JS bridge
    // expects (matching shader.customBindings).
    std::vector<std::uint32_t> handles;
    handles.reserve(shader.customBindings.size());
    const std::uint8_t* p = static_cast<const std::uint8_t*>(push);
    for (const auto& b : shader.customBindings) {
        if (b.pushOffset + sizeof(std::uint32_t) > pushBytes) {
            std::println("UIRenderer::Dispatch: binding pushOffset {} out of bounds (push={})",
                         b.pushOffset, pushBytes);
            return;
        }
        std::uint32_t slot;
        std::memcpy(&slot, p + b.pushOffset, sizeof(slot));
        std::uint32_t handle = 0;
        switch (b.kind) {
            case UICustomBindingKind::Buffer:
                if (slot < heap_->bufferTable.size()) handle = heap_->bufferTable[slot];
                break;
            case UICustomBindingKind::SampledTexture:
                if (slot < heap_->imageTable.size()) handle = heap_->imageTable[slot];
                break;
            case UICustomBindingKind::Sampler:
                if (slot < heap_->samplerTable.size()) handle = heap_->samplerTable[slot];
                break;
        }
        handles.push_back(handle);
    }
    WebGPU::wgpuDispatchCustom(shader.pipelineHandle,
                               push, static_cast<std::int32_t>(pushBytes),
                               handles.data(), static_cast<std::int32_t>(handles.size()),
                               static_cast<std::int32_t>(gx),
                               static_cast<std::int32_t>(gy),
                               static_cast<std::int32_t>(gz));
 }
--- a/implementations/Crafter.Graphics-UI.cpp
+++ b/implementations/Crafter.Graphics-UI.cpp
@ -28,14 +28,13 @@ import :ImageVulkan;
 import :VulkanBuffer;
 import :FontAtlas;
 import :Font;
 import :GraphicsTypes;
 import std;
 using namespace Crafter;
 // ─── Initialize ─────────────────────────────────────────────────────────
-void UIRenderer::Initialize(Window& window, GraphicsDescriptorHeap& heap, GraphicsCommandBuffer initCmd,
+void UIRenderer::Initialize(Window& window, DescriptorHeapVulkan& heap, VkCommandBuffer initCmd,
                            std::filesystem::path quadsSpv,
                            std::filesystem::path circlesSpv,
                            std::filesystem::path imagesSpv,
@ -82,7 +81,7 @@ void UIRenderer::Initialize(Window& window, GraphicsDescriptorHeap& heap, Graphi
 // ─── per-frame Record ───────────────────────────────────────────────────
-void UIRenderer::Record(GraphicsCommandBuffer cmd, std::uint32_t frameIdx, Window& window) {
+void UIRenderer::Record(VkCommandBuffer cmd, std::uint32_t frameIdx, Window& window) {
    // Reset per-frame state.
    firstDispatchThisFrame_ = true;
@ -103,6 +102,28 @@ void UIRenderer::Record(GraphicsCommandBuffer cmd, std::uint32_t frameIdx, Windo
    (void)window;
 }
 // ─── header builder ─────────────────────────────────────────────────────
 UIDispatchHeader UIRenderer::FillHeader(std::uint32_t itemBufferSlot,
                                        std::uint32_t itemCount,
                                        std::array<float,4> clipRectPx,
                                        std::uint32_t flags) const noexcept {
    UIDispatchHeader h{};
    h.outImage      = outImageSlot_;
    h.itemBuffer    = itemBufferSlot;
    h.surfaceWidth  = window_->width;
    h.surfaceHeight = window_->height;
    h.clipX = clipRectPx[0];
    h.clipY = clipRectPx[1];
    h.clipW = clipRectPx[2];
    h.clipH = clipRectPx[3];
    h.itemCount = itemCount;
    h.frameIdx  = window_->currentBuffer;
    h.flags     = flags;
    h._pad      = 0;
    return h;
 }
 // ─── group-count helper ────────────────────────────────────────────────
 namespace {
@ -121,7 +142,7 @@ namespace {
 // without inter-workgroup races on imageLoad/imageStore — the bug that the
 // per-item dispatch model had.
-void UIRenderer::DispatchQuads(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot,
+void UIRenderer::DispatchQuads(VkCommandBuffer cmd, std::uint32_t bufferSlot,
                               std::uint32_t itemCount,
                               std::array<float,4> clipRectPx) {
    if (itemCount == 0) return;
@ -130,7 +151,7 @@ void UIRenderer::DispatchQuads(GraphicsCommandBuffer cmd, std::uint32_t bufferSl
             TilesFor(window_->width), TilesFor(window_->height), 1u);
 }
-void UIRenderer::DispatchCircles(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot,
+void UIRenderer::DispatchCircles(VkCommandBuffer cmd, std::uint32_t bufferSlot,
                                 std::uint32_t itemCount,
                                 std::array<float,4> clipRectPx) {
    if (itemCount == 0) return;
@ -139,7 +160,7 @@ void UIRenderer::DispatchCircles(GraphicsCommandBuffer cmd, std::uint32_t buffer
             TilesFor(window_->width), TilesFor(window_->height), 1u);
 }
-void UIRenderer::DispatchImages(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot,
+void UIRenderer::DispatchImages(VkCommandBuffer cmd, std::uint32_t bufferSlot,
                                std::uint32_t itemCount,
                                std::array<float,4> clipRectPx) {
    if (itemCount == 0) return;
@ -148,7 +169,7 @@ void UIRenderer::DispatchImages(GraphicsCommandBuffer cmd, std::uint32_t bufferS
             TilesFor(window_->width), TilesFor(window_->height), 1u);
 }
-void UIRenderer::DispatchText(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot,
+void UIRenderer::DispatchText(VkCommandBuffer cmd, std::uint32_t bufferSlot,
                              std::uint32_t itemCount,
                              std::array<float,4> clipRectPx) {
    if (itemCount == 0) return;
@ -178,7 +199,7 @@ void UIRenderer::DispatchText(GraphicsCommandBuffer cmd, std::uint32_t bufferSlo
 // ─── generic Dispatch (with barrier) ────────────────────────────────────
-void UIRenderer::Dispatch(GraphicsCommandBuffer cmd, const GraphicsComputeShader& shader,
+void UIRenderer::Dispatch(VkCommandBuffer cmd, const ComputeShader& shader,
                          const void* push, std::uint32_t pushBytes,
                          std::uint32_t gx, std::uint32_t gy, std::uint32_t gz) {
    if (!firstDispatchThisFrame_) {
@ -335,3 +356,46 @@ SamplerSlot UIRenderer::RegisterLinearClampSampler() {
    return RegisterSampler(s);
 }
 // ─── ShapeText ─────────────────────────────────────────────────────────
 std::uint32_t UIRenderer::ShapeText(Font& font, float pxSize,
                                    float x, float baselineY,
                                    std::string_view utf8,
                                    std::array<float,4> color,
                                    GlyphItem* out, std::uint32_t outCapacity,
                                    float* outAdvance) {
    if (fontAtlas == nullptr) {
        throw std::runtime_error("UIRenderer::ShapeText: no FontAtlas (set fontAtlas before Initialize)");
    }
    const float scale = pxSize / FontAtlas::kBaseSize;
    float cursor = x;
    std::uint32_t written = 0;
    std::size_t i = 0;
    while (i < utf8.size() && written < outCapacity) {
        std::uint32_t cp = DecodeUtf8(utf8, i);
        if (cp == 0) break;
        if (cp == '\n') { /* single-line shaper — ignore */ continue; }
        fontAtlas->Ensure(font, cp);
        const Glyph* g = fontAtlas->Lookup(font, cp);
        if (g == nullptr) continue;
        // Empty glyph (whitespace) — advance only.
        if (g->w > 0 && g->h > 0) {
            GlyphItem& gi = out[written++];
            gi.x  = cursor + g->xoff * scale;
            gi.y  = baselineY + g->yoff * scale;
            gi.w  = g->w * scale;
            gi.h  = g->h * scale;
            gi.u0 = g->u0; gi.v0 = g->v0;
            gi.u1 = g->u1; gi.v1 = g->v1;
            gi.r = color[0]; gi.g = color[1]; gi.b = color[2]; gi.a = color[3];
        }
        cursor += g->advance * scale;
    }
    if (outAdvance) *outAdvance = cursor - x;
    return written;
 }
--- a/implementations/Crafter.Graphics-WebGPUComputeShader.cpp
+++ b/implementations/Crafter.Graphics-WebGPUComputeShader.cpp
@ -1,41 +0,0 @@
 /*
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 module Crafter.Graphics:WebGPUComputeShader_impl;
 import :WebGPUComputeShader;
 import :WebGPU;
 import std;
 using namespace Crafter;
 void WebGPUComputeShader::Load(std::string_view wgsl,
                               std::span<const UICustomBinding> bindings) {
    customBindings.assign(bindings.begin(), bindings.end());
    pipelineHandle = WebGPU::wgpuLoadCustomShader(
        wgsl.data(),
        static_cast<std::int32_t>(wgsl.size()),
        customBindings.data(),
        static_cast<std::int32_t>(customBindings.size())
    );
 }
 void WebGPUComputeShader::Load(const std::filesystem::path& wgslPath,
                               std::span<const UICustomBinding> bindings) {
    std::ifstream f(wgslPath, std::ios::binary | std::ios::ate);
    if (!f.is_open()) {
        std::println("WebGPUComputeShader::Load: cannot open {}", wgslPath.string());
        return;
    }
    auto size = f.tellg();
    if (size <= 0) {
        std::println("WebGPUComputeShader::Load: empty file {}", wgslPath.string());
        return;
    }
    f.seekg(0, std::ios::beg);
    std::string src(static_cast<std::size_t>(size), '\0');
    f.read(src.data(), size);
    Load(std::string_view{src}, bindings);
 }
--- a/implementations/Crafter.Graphics-Window.cpp
+++ b/implementations/Crafter.Graphics-Window.cpp
@ -69,10 +69,6 @@ import :Gamepad;
 import :VulkanTransition;
 import :DescriptorHeapVulkan;
 import :RenderPass;
 #ifdef CRAFTER_GRAPHICS_WINDOW_DOM
 import :WebGPU;
 import :DescriptorHeapWebGPU;
 #endif
 import std;
 using namespace Crafter;
@ -1327,22 +1323,23 @@ Window::Window(std::uint32_t w, std::uint32_t h, const std::string_view title)
 Window::Window(std::uint32_t w, std::uint32_t h) : width(w), height(h) {
    if (g_domWindow != nullptr) {
        // Only one Window per page in V1. Subsequent constructions are
        // a programming error — log loudly and clobber the previous
        // pointer so the new Window's events at least fire.
        // (stderr isn't reachable via `import std;` on wasi-sdk yet; just log
        // to cout. The browser console pipes both to the same place.)
        std::println("Crafter::Window: only one DOM Window per page; "
                     "overwriting the previous instance.");
    }
    g_domWindow = this;
-    // Browser owns the real surface size. The width/height passed in are
+    // Use the browser-reported viewport size as the initial dimensions
-    // advisory only — useful as a native-side hint, ignored on DOM. We
+    // unless the caller asked for something specific. Browser owns the
-    // always sync to innerWidth/innerHeight so:
+    // real size; w/h passed in are advisory.
-    //   - window.width/.height match the canvas's CSS pixel size,
+    if (w == 0 || h == 0) {
    //   - MouseEvent.clientX/.clientY (CSS pixels) compare correctly
    //     against any layout done with window.width/.height,
    //   - the dispatch group count from window.width/8 covers the
    //     canvas exactly.
    (void)w; (void)h;
        width  = static_cast<std::uint32_t>(Crafter::DomEnv::domGetInnerWidth());
        height = static_cast<std::uint32_t>(Crafter::DomEnv::domGetInnerHeight());
    }
    // The handle passed to attach is just a non-zero token the JS side
    // includes back in every dispatcher call. We don't use it on the
@ -1397,18 +1394,10 @@ void Window::SetDefaultCursor() {
 }
 void Window::StartSync() {
-    // Hand the loop to rAF, then exit the wasm via _Exit so wasi-libc
+    // Hand the loop to rAF. Returns immediately; the wasm `_start`
-    // skips __wasm_call_dtors. If we let main return normally, _start
+    // (main) finishes, and the runtime keeps the module alive while
-    // calls __wasm_call_dtors → static destructors fire (including
+    // the JS-side rAF chain ticks `__crafterDom_frame`.
    // Window's own), then __wasi_proc_exit → wasm trap. Subsequent rAF
    // calls into the wasm would then trap too, killing rendering.
    // _Exit jumps straight to __wasi_proc_exit, which our runtime.js
    // catches via a thrown sentinel so the instance stays alive while
    // every static-allocated object (Window, UIRenderer, GPU buffers,
    // event listeners) remains untouched. Callers' code after
    // StartSync() never runs — match that contract on native too.
    Crafter::DomEnv::domStartFrameLoop();
    std::_Exit(0);
 }
 void Window::StartUpdate() {
@ -1432,25 +1421,10 @@ void Window::Update() {
 }
 void Window::Render() {
-    if (!open) return;
+    // V1: no rendering in DOM mode. Kept as a callable no-op so
-    Crafter::WebGPU::wgpuFrameBegin();
+    // existing cross-platform code paths (e.g. main loops calling
-    for (RenderPass* p : passes) {
+    // window.Render() before window.StartSync()) compile. V2 will
-        if (p) p->Record(/*cmd*/ 0u, currentBuffer, *this);
+    // hang the WebGPU command-submit here.
    }
    Crafter::WebGPU::wgpuFrameEnd();
 }
 WebGPUCommandEncoderRef Window::StartInit() {
    // DOM init: no command buffer needed — texture / buffer creation goes
    // through synchronous wgpu* imports. Return 0 as a placeholder; the
    // value is opaque to user code (auto-typed in HelloUI).
    Crafter::WebGPU::wgpuInit();
    return 0;
 }
 void Window::FinishInit() {
    // Nothing to submit in DOM mode; all init writes are queued at call
    // time via queue.writeBuffer / writeTexture.
 }
 // ─── C exports the JS bridge calls back into ──────────────────────────
@ -1463,7 +1437,6 @@ extern "C" {
        g_domWindow->onBeforeUpdate.Invoke();
        if (g_domWindow->updating) {
            g_domWindow->Update();
            g_domWindow->Render();
        }
    }
--- a/interfaces/Crafter.Graphics-DescriptorHeapWebGPU.cppm
+++ b/interfaces/Crafter.Graphics-DescriptorHeapWebGPU.cppm
@ -1,185 +0,0 @@
 /*
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 // DOM-mode parallel to DescriptorHeapVulkan. WebGPU has no real bindless,
 // so the "heap" is purely a CPU-side slot allocator with a side-table
 // mapping slot → JS-side WebGPU handle. UIRenderer looks up the handle by
 // slot at dispatch time to build (or fetch from cache) the bind group.
 export module Crafter.Graphics:DescriptorHeapWebGPU;
 #ifdef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 import :WebGPU;
 export namespace Crafter {
    struct DescriptorHeapWebGPU;
    struct DescriptorRange {
        std::uint16_t firstElement;
        std::uint16_t count;
    };
    class BufferSlot {
    public:
        DescriptorHeapWebGPU* heap = nullptr;
        std::uint16_t firstElement = 0xFFFF;
        BufferSlot() = default;
        BufferSlot(DescriptorHeapWebGPU* h, std::uint16_t f) : heap(h), firstElement(f) {}
        BufferSlot(const BufferSlot&) = delete;
        BufferSlot& operator=(const BufferSlot&) = delete;
        BufferSlot(BufferSlot&& o) noexcept : heap(o.heap), firstElement(o.firstElement) { o.firstElement = 0xFFFF; }
        BufferSlot& operator=(BufferSlot&& o) noexcept;
        ~BufferSlot();
        explicit operator bool() const noexcept { return firstElement != 0xFFFF; }
        operator std::uint16_t() const noexcept { return firstElement; }
    };
    class ImageSlot {
    public:
        DescriptorHeapWebGPU* heap = nullptr;
        std::uint16_t firstElement = 0xFFFF;
        ImageSlot() = default;
        ImageSlot(DescriptorHeapWebGPU* h, std::uint16_t f) : heap(h), firstElement(f) {}
        ImageSlot(const ImageSlot&) = delete;
        ImageSlot& operator=(const ImageSlot&) = delete;
        ImageSlot(ImageSlot&& o) noexcept : heap(o.heap), firstElement(o.firstElement) { o.firstElement = 0xFFFF; }
        ImageSlot& operator=(ImageSlot&& o) noexcept;
        ~ImageSlot();
        explicit operator bool() const noexcept { return firstElement != 0xFFFF; }
        operator std::uint16_t() const noexcept { return firstElement; }
    };
    class SamplerSlot {
    public:
        DescriptorHeapWebGPU* heap = nullptr;
        std::uint16_t firstElement = 0xFFFF;
        SamplerSlot() = default;
        SamplerSlot(DescriptorHeapWebGPU* h, std::uint16_t f) : heap(h), firstElement(f) {}
        SamplerSlot(const SamplerSlot&) = delete;
        SamplerSlot& operator=(const SamplerSlot&) = delete;
        SamplerSlot(SamplerSlot&& o) noexcept : heap(o.heap), firstElement(o.firstElement) { o.firstElement = 0xFFFF; }
        SamplerSlot& operator=(SamplerSlot&& o) noexcept;
        ~SamplerSlot();
        explicit operator bool() const noexcept { return firstElement != 0xFFFF; }
        operator std::uint16_t() const noexcept { return firstElement; }
    };
    struct DescriptorHeapWebGPU {
        std::vector<WebGPUBufferRef>  bufferTable;
        std::vector<WebGPUTextureRef> imageTable;
        std::vector<WebGPUSamplerRef> samplerTable;
        std::vector<std::uint16_t> bufferFreelist;
        std::vector<std::uint16_t> imageFreelist;
        std::vector<std::uint16_t> samplerFreelist;
        std::uint16_t nextBuffer  = 0;
        std::uint16_t nextImage   = 0;
        std::uint16_t nextSampler = 0;
        void Initialize(std::uint16_t images, std::uint16_t buffers, std::uint16_t samplers) {
            imageTable.assign(images, 0);
            bufferTable.assign(buffers, 0);
            samplerTable.assign(samplers, 0);
            imageFreelist.reserve(images);
            bufferFreelist.reserve(buffers);
            samplerFreelist.reserve(samplers);
        }
        DescriptorRange AllocateBufferSlots(std::uint16_t count) {
            if (count == 1 && !bufferFreelist.empty()) {
                auto f = bufferFreelist.back(); bufferFreelist.pop_back();
                return { f, 1 };
            }
            if (nextBuffer + count > bufferTable.size()) {
                std::println("DescriptorHeapWebGPU: buffer slots exhausted");
                std::abort();
            }
            DescriptorRange r{ nextBuffer, count };
            nextBuffer = static_cast<std::uint16_t>(nextBuffer + count);
            return r;
        }
        DescriptorRange AllocateImageSlots(std::uint16_t count) {
            if (count == 1 && !imageFreelist.empty()) {
                auto f = imageFreelist.back(); imageFreelist.pop_back();
                return { f, 1 };
            }
            if (nextImage + count > imageTable.size()) {
                std::println("DescriptorHeapWebGPU: image slots exhausted");
                std::abort();
            }
            DescriptorRange r{ nextImage, count };
            nextImage = static_cast<std::uint16_t>(nextImage + count);
            return r;
        }
        DescriptorRange AllocateSamplerSlots(std::uint16_t count) {
            if (count == 1 && !samplerFreelist.empty()) {
                auto f = samplerFreelist.back(); samplerFreelist.pop_back();
                return { f, 1 };
            }
            if (nextSampler + count > samplerTable.size()) {
                std::println("DescriptorHeapWebGPU: sampler slots exhausted");
                std::abort();
            }
            DescriptorRange r{ nextSampler, count };
            nextSampler = static_cast<std::uint16_t>(nextSampler + count);
            return r;
        }
        void FreeBufferSlots(std::uint16_t first, std::uint16_t count) noexcept {
            for (std::uint16_t i = 0; i < count; ++i) {
                bufferTable[first + i] = 0;
                bufferFreelist.push_back(first + i);
            }
        }
        void FreeImageSlots(std::uint16_t first, std::uint16_t count) noexcept {
            for (std::uint16_t i = 0; i < count; ++i) {
                imageTable[first + i] = 0;
                imageFreelist.push_back(first + i);
            }
        }
        void FreeSamplerSlots(std::uint16_t first, std::uint16_t count) noexcept {
            for (std::uint16_t i = 0; i < count; ++i) {
                samplerTable[first + i] = 0;
                samplerFreelist.push_back(first + i);
            }
        }
    };
    // ─── slot dtors (defined here since they reference DescriptorHeapWebGPU) ─
    inline BufferSlot::~BufferSlot() {
        if (firstElement != 0xFFFF && heap) heap->FreeBufferSlots(firstElement, 1);
    }
    inline BufferSlot& BufferSlot::operator=(BufferSlot&& o) noexcept {
        if (this != &o) {
            if (firstElement != 0xFFFF && heap) heap->FreeBufferSlots(firstElement, 1);
            heap = o.heap; firstElement = o.firstElement; o.firstElement = 0xFFFF;
        }
        return *this;
    }
    inline ImageSlot::~ImageSlot() {
        if (firstElement != 0xFFFF && heap) heap->FreeImageSlots(firstElement, 1);
    }
    inline ImageSlot& ImageSlot::operator=(ImageSlot&& o) noexcept {
        if (this != &o) {
            if (firstElement != 0xFFFF && heap) heap->FreeImageSlots(firstElement, 1);
            heap = o.heap; firstElement = o.firstElement; o.firstElement = 0xFFFF;
        }
        return *this;
    }
    inline SamplerSlot::~SamplerSlot() {
        if (firstElement != 0xFFFF && heap) heap->FreeSamplerSlots(firstElement, 1);
    }
    inline SamplerSlot& SamplerSlot::operator=(SamplerSlot&& o) noexcept {
        if (this != &o) {
            if (firstElement != 0xFFFF && heap) heap->FreeSamplerSlots(firstElement, 1);
            heap = o.heap; firstElement = o.firstElement; o.firstElement = 0xFFFF;
        }
        return *this;
    }
 }
 #endif // CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-Font.cppm
+++ b/interfaces/Crafter.Graphics-Font.cppm
@ -19,8 +19,13 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301 USA
 */
 module;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "../lib/stb_truetype.h"
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
 export module Crafter.Graphics:Font;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 namespace Crafter {
@ -31,6 +36,7 @@ namespace Crafter {
        if (i >= text.size()) return 0;
        std::uint8_t b0 = static_cast<std::uint8_t>(text[i]);
        // Single-byte ASCII is the common path.
        if (b0 < 0x80) { ++i; return b0; }
        int extra;
@ -38,13 +44,13 @@ namespace Crafter {
        if      ((b0 & 0xE0) == 0xC0) { extra = 1; cp = b0 & 0x1F; }
        else if ((b0 & 0xF0) == 0xE0) { extra = 2; cp = b0 & 0x0F; }
        else if ((b0 & 0xF8) == 0xF0) { extra = 3; cp = b0 & 0x07; }
-        else { ++i; return 0xFFFD; }
+        else { ++i; return 0xFFFD; }   // continuation byte at start, or 5+-byte leader
        ++i;
        for (int k = 0; k < extra; ++k) {
            if (i >= text.size()) return 0xFFFD;
            std::uint8_t b = static_cast<std::uint8_t>(text[i]);
-            if ((b & 0xC0) != 0x80) return 0xFFFD;
+            if ((b & 0xC0) != 0x80) return 0xFFFD;       // missing continuation
            cp = (cp << 6) | (b & 0x3Fu);
            ++i;
        }
@ -61,7 +67,8 @@ namespace Crafter {
        Font(const std::filesystem::path& font);
        std::uint32_t GetLineWidth(const std::string_view text, float size);
        float LineHeight(float size);
-        float AscentPx(float size);
+        float AscentPx(float size);     // baseline offset from line-top
-        float ScaleForSize(float size);
+        float ScaleForSize(float size); // stb's pixel-units-per-em factor
    };
 }
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-FontAtlas.cppm
+++ b/interfaces/Crafter.Graphics-FontAtlas.cppm
@ -19,61 +19,70 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301 USA
 module;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
-#endif
+#endif // !CRAFTER_GRAPHICS_WINDOW_DOM
 export module Crafter.Graphics:FontAtlas;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 import :Font;
 import :GraphicsTypes;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import :ImageVulkan;
 import :Device;
 #else
 import :WebGPU;
 #endif
 export namespace Crafter {
    // Per-glyph metrics. UVs are 0..1 in atlas space; on-screen sizes /
    // offsets / advance are in *atlas pixels at the base size* and scale
    // linearly with the requested font size at draw time.
    struct Glyph {
-        float u0 = 0, v0 = 0;
+        float u0 = 0, v0 = 0;         // top-left UV in the atlas
-        float u1 = 0, v1 = 0;
+        float u1 = 0, v1 = 0;         // bottom-right UV in the atlas
-        float w  = 0, h  = 0;
+        float w  = 0, h  = 0;         // glyph quad size in atlas px (= the bitmap size)
-        float xoff = 0, yoff = 0;
+        float xoff = 0, yoff = 0;     // glyph bearing relative to baseline cursor
-        float advance = 0;
+        float advance = 0;            // horizontal advance at base size, in atlas px
    };
    // Single-channel SDF atlas. Glyphs are rasterised with stb_truetype's
    // GetGlyphSDF at a fixed `kBaseSize` resolution and packed via a simple
    // shelf allocator. Drawing scales the glyph quad linearly; the shader
    // resolves edge AA via screen-space derivatives, so a single atlas
    // serves all sizes and DPI scales without re-bake.
    class FontAtlas {
    public:
        // Build-time constants. Tweak in one place if needed; values picked
        // to give crisp text from ~10pt to ~96pt and leave headroom in the
        // SDF distance band so smoothstep is in the linear regime.
        static constexpr int   kAtlasSize     = 1024;
-        static constexpr float kBaseSize      = 32.0f;
+        static constexpr float kBaseSize      = 32.0f;     // pixel-height at which we rasterise
-        static constexpr int   kPadding       = 4;
+        static constexpr int   kPadding       = 4;         // distance-field padding around each glyph
-        static constexpr int   kOnEdgeValue   = 128;
+        static constexpr int   kOnEdgeValue   = 128;       // 8-bit value mapped to "0 distance"
-        static constexpr float kPixelDistScale = 32.0f;
+        static constexpr float kPixelDistScale = 32.0f;    // how many distance units per pixel — wider = softer AA range
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
        ImageVulkan<std::uint8_t> image;
-#else
+        bool dirty = false;        // staging has unflushed writes
        WebGPUTextureRef textureHandle = 0;
        std::vector<std::uint8_t> staging;
 #endif
        bool dirty = false;
-        void Initialize(GraphicsCommandBuffer cmd);
+        // Allocate the GPU image and zero-clear it. Must be called once
-
+        // with a one-shot init command buffer.
-        // Returns the row-major byte pointer the CPU writes pixels into.
+        void Initialize(VkCommandBuffer cmd);
        // Same shape on both backends.
        std::uint8_t* PixelPtr() noexcept;
        // Rasterise + pack the glyph if it isn't cached yet. Returns
        // false only if the atlas is out of space (V2: grow). After a
        // successful Ensure the bitmap lives in `image.buffer.value` and
        // `dirty` is true; call Update(cmd) before reading on the GPU.
        bool Ensure(Font& font, std::uint32_t codepoint);
        // Lookup is cheap (hash-table). Returns nullptr if the glyph
        // hasn't been Ensured.
        const Glyph* Lookup(Font& font, std::uint32_t codepoint) const;
-        void Update(GraphicsCommandBuffer cmd);
+
        // If `dirty`, flushes staging into the GPU image and transitions
        // it back to SHADER_READ_ONLY_OPTIMAL. No-op if not dirty.
        void Update(VkCommandBuffer cmd);
    private:
        // Shelf packer state.
        struct Shelf { int y = 0; int height = 0; int cursorX = 0; };
        std::vector<Shelf> shelves_;
        int  nextShelfY_ = 0;
        // (font*, codepoint) → Glyph cache.
        struct Key {
            const Font* font;
            std::uint32_t cp;
@ -88,6 +97,8 @@ export namespace Crafter {
        };
        std::unordered_map<Key, Glyph, KeyHash> cache_;
        // Place a wxh glyph; returns true + writes top-left into outX/outY.
        bool ShelfPlace(int w, int h, int& outX, int& outY);
    };
 }
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-GraphicsTypes.cppm
+++ b/interfaces/Crafter.Graphics-GraphicsTypes.cppm
@ -1,42 +0,0 @@
 /*
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 // Backend-portable type aliases. NOT an abstraction layer — these are pure
 // `using` declarations that resolve to the backend's native types per the
 // active CRAFTER_GRAPHICS_WINDOW_* define.
 module;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
 #endif
 export module Crafter.Graphics:GraphicsTypes;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import :VulkanBuffer;
 import :DescriptorHeapVulkan;
 import :ComputeShader;
 export namespace Crafter {
    using GraphicsCommandBuffer = VkCommandBuffer;
    using GraphicsDescriptorHeap = DescriptorHeapVulkan;
    using GraphicsComputeShader = ComputeShader;
    template<class T, bool Mapped>
    using GraphicsBuffer = VulkanBuffer<T, Mapped>;
 }
 #else
 import :WebGPU;
 import :WebGPUBuffer;
 import :DescriptorHeapWebGPU;
 import :WebGPUComputeShader;
 export namespace Crafter {
    using GraphicsCommandBuffer = WebGPUCommandEncoderRef;
    using GraphicsDescriptorHeap = DescriptorHeapWebGPU;
    using GraphicsComputeShader = WebGPUComputeShader;
    template<class T, bool Mapped>
    using GraphicsBuffer = WebGPUBuffer<T, Mapped>;
 }
 #endif
--- a/interfaces/Crafter.Graphics-RenderPass.cppm
+++ b/interfaces/Crafter.Graphics-RenderPass.cppm
@ -16,15 +16,20 @@ You should have received a copy of the GNU Lesser General Public
 License along with this library; if not, write to the Free Software
 Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301 USA
 */
 module;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
 export module Crafter.Graphics:RenderPass;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 import :GraphicsTypes;
 export namespace Crafter {
    struct Window;
    struct RenderPass {
-        virtual void Record(GraphicsCommandBuffer cmd, std::uint32_t frameIdx, Window& window) = 0;
+        virtual void Record(VkCommandBuffer cmd, std::uint32_t frameIdx, Window& window) = 0;
        virtual ~RenderPass() = default;
    };
 }
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-UI.cppm
+++ b/interfaces/Crafter.Graphics-UI.cppm
@ -19,27 +19,20 @@ Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301 USA
 module;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
-#endif
+#endif // !CRAFTER_GRAPHICS_WINDOW_DOM
 export module Crafter.Graphics:UI;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 import Crafter.Event;
 import :Device;
 import :Window;
 import :RenderPass;
 import :GraphicsTypes;
 import :FontAtlas;
 import :Font;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import :Device;
 import :DescriptorHeapVulkan;
 import :ImageVulkan;
 import :VulkanBuffer;
 import :ComputeShader;
-#else
+import :FontAtlas;
-import :DescriptorHeapWebGPU;
+import :Font;
 import :WebGPU;
 import :WebGPUBuffer;
 import :WebGPUComputeShader;
 #endif
 export namespace Crafter {
    // ─── push-constant header ───────────────────────────────────────────
@ -63,8 +56,8 @@ export namespace Crafter {
    struct QuadItem {
        float x, y, w, h;
        float r, g, b, a;
-        float cTL, cTR, cBR, cBL;
+        float cTL, cTR, cBR, cBL;       // per-corner radius in px
-        float outline, oR, oG, oB;
+        float outline, oR, oG, oB;      // outline thickness + RGB
    };
    static_assert(sizeof(QuadItem) == 64);
@ -91,11 +84,15 @@ export namespace Crafter {
    static_assert(sizeof(GlyphItem) == 48);
    // ─── tiny rect-carving helper ───────────────────────────────────────
    // Pure value semantics. No engine, just convenience. Skip if you'd rather
    // compute pixels yourself.
    struct Rect {
        float x = 0, y = 0, w = 0, h = 0;
        enum class Anchor { Top, Bottom, Left, Right };
        // Returns a sub-rect of `size` along the given anchor edge of self.
        // Does not modify `*this`. (Use `.Inset(...)` to drop a margin first.)
        Rect SubRect(float size, Anchor a) const noexcept {
            switch (a) {
                case Anchor::Top:    return { x, y,                 w,    std::min(size, h) };
@ -126,74 +123,121 @@ export namespace Crafter {
    // ─── per-frame callback args ────────────────────────────────────────
    struct UIBuildArgs {
-        GraphicsCommandBuffer cmd;
+        VkCommandBuffer cmd;
        std::uint32_t   frameIdx;
    };
    // ─── UIRenderer ─────────────────────────────────────────────────────
    // One per Window (typically). Owns the four standard compute shaders,
    // pre-allocates heap slots for the swapchain images, and exposes a thin
    // dispatch helper for both the standard shaders and user-supplied ones.
    //
    // Workflow:
    //   1. Construct, configure (set fontAtlas if drawing text).
    //   2. Initialize(window, heap, initCmd) — once, after window.descriptorHeap
    //      is set and before window.FinishInit().
    //   3. window.passes.push_back(&ui).
    //   4. Listen on `onBuild`. Inside the callback, fill your item buffers,
    //      flush them, and call DispatchQuads / DispatchCircles / DispatchImages
    //      / DispatchText / Dispatch as needed. Library inserts a SHADER_WRITE
    //      → SHADER_READ|WRITE memory barrier between consecutive dispatches.
    class UIRenderer : public RenderPass {
    public:
-        GraphicsComputeShader drawQuads;
+        // Pre-loaded standard shaders (public so users can call Dispatch
-        GraphicsComputeShader drawCircles;
+        // directly with them if they want to embed extra push-constant fields
-        GraphicsComputeShader drawImages;
+        // beyond the standard header).
-        GraphicsComputeShader drawText;
+        ComputeShader drawQuads;
        ComputeShader drawCircles;
        ComputeShader drawImages;
        ComputeShader drawText;
        // Optional. If set before Initialize, the atlas is registered into a
        // sampled-image slot + linear sampler slot, and Update(cmd) is called
        // at the top of every Record() so any glyphs ensured during onBuild
        // make it to the GPU before the text dispatch reads them.
        FontAtlas* fontAtlas = nullptr;
        // User callback. Subscribe by holding a Crafter::EventListener<UIBuildArgs>:
        //     EventListener<UIBuildArgs> sub(&ui.onBuild, [&](UIBuildArgs a) { ... });
        // Listener lifetime governs the subscription.
        Crafter::Event<UIBuildArgs> onBuild;
        UIRenderer() = default;
        UIRenderer(const UIRenderer&) = delete;
        UIRenderer& operator=(const UIRenderer&) = delete;
-        void Initialize(Window& window, GraphicsDescriptorHeap& heap, GraphicsCommandBuffer initCmd,
+        // Default shader paths assume Crafter.Build placed the .spv files
        // alongside the consumer binary (this is what cfg.shaders does).
        void Initialize(Window& window, DescriptorHeapVulkan& heap, VkCommandBuffer initCmd,
                        std::filesystem::path quadsSpv   = "ui-quads.comp.spv",
                        std::filesystem::path circlesSpv = "ui-circles.comp.spv",
                        std::filesystem::path imagesSpv  = "ui-images.comp.spv",
                        std::filesystem::path textSpv    = "ui-text.comp.spv");
-        void Record(GraphicsCommandBuffer cmd, std::uint32_t frameIdx, Window& window) override;
+        // RenderPass interface — invoked from Window::Render.
        void Record(VkCommandBuffer cmd, std::uint32_t frameIdx, Window& window) override;
        // ─── helpers used inside `onBuild` ─────────────────────────────
        // Builds a populated header. `clipRectPx` defaults to "no clip".
        UIDispatchHeader FillHeader(std::uint32_t itemBufferSlot,
                                    std::uint32_t itemCount,
                                    std::array<float,4> clipRectPx = {0.0f, 0.0f, 1e9f, 1e9f},
                                    std::uint32_t flags = 0) const noexcept;
-        void DispatchQuads(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
+        // Convenience: dispatches the named standard shader. Group count is
        // computed from the window's surface size — the standard shaders
        // dispatch one workgroup per 8×8 screen tile and iterate every item
        // in the buffer in order, so item ORDER in the buffer == draw order
        // on screen (later items overdraw earlier ones, race-free).
        void DispatchQuads(VkCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
                           std::array<float,4> clipRectPx = {0.0f, 0.0f, 1e9f, 1e9f});
-        void DispatchCircles(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
+        void DispatchCircles(VkCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
                             std::array<float,4> clipRectPx = {0.0f, 0.0f, 1e9f, 1e9f});
-        void DispatchImages(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
+        void DispatchImages(VkCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
                            std::array<float,4> clipRectPx = {0.0f, 0.0f, 1e9f, 1e9f});
-        void DispatchText(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
+        // For DispatchText, the font atlas image+sampler slots are taken from
        // UIRenderer's Initialize-time registration (see fontAtlasImageSlot()
        // / fontAtlasSamplerSlot()). Set `fontAtlas` before Initialize.
        void DispatchText(VkCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
                          std::array<float,4> clipRectPx = {0.0f, 0.0f, 1e9f, 1e9f});
-        // Generic dispatch for user-authored shaders. On Vulkan, `shader` is
+        // Generic dispatch — for user-authored shaders. Inserts the standard
-        // a SPIR-V compute pipeline (bindless via VK_EXT_descriptor_heap, so
+        // pre-dispatch barrier (skipped on the first call per frame).
-        // any resource indices baked into push data resolve through the
+        void Dispatch(VkCommandBuffer cmd, const ComputeShader& shader,
        // global heap). On DOM, `shader` carries a UICustomBinding list
        // declared at Load time; the renderer reads the listed slot uints
        // out of `push`, resolves them against heap.bufferTable /
        // imageTable / samplerTable, and builds the bind groups before
        // dispatching.
        void Dispatch(GraphicsCommandBuffer cmd, const GraphicsComputeShader& shader,
                      const void* push, std::uint32_t pushBytes,
                      std::uint32_t gx, std::uint32_t gy = 1, std::uint32_t gz = 1);
-        // Allocates a heap slot for the buffer and registers the GPU handle.
+        // Allocates a heap slot for the buffer and writes its descriptor into
-        // Returns a move-only BufferSlot RAII handle.
+        // every per-frame heap. The user's mapped buffer is shared across
        // frames — fine because Window::Render currently waits idle before
        // submitting the next frame. Returns a move-only BufferSlot handle
        // whose destructor returns the slot to the heap. Implicitly converts
        // to the absolute heap index when passed to FillHeader / Dispatch*.
        template<typename T, bool Mapped>
-        BufferSlot RegisterBuffer(GraphicsBuffer<T, Mapped>& buffer);
+        BufferSlot RegisterBuffer(VulkanBuffer<T, Mapped>& buffer);
-#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
+        // Same for an ImageVulkan-managed sampled image (e.g. a user texture).
        // Caller specifies the layout the image will be sampled in (typically
        // VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL).
        template<typename Pixel>
        ImageSlot RegisterImage(ImageVulkan<Pixel>& image, VkFormat format,
                                VkImageLayout layout = VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL);
        // Allocates a sampler slot and writes a VkSamplerCreateInfo into
        // every per-frame sampler heap. v1 takes the create-info inline.
        SamplerSlot RegisterSampler(const VkSamplerCreateInfo& info);
-#endif
+
        // Convenience: a linear-filter, clamp-to-edge sampler. Returns a
        // SamplerSlot handle. Useful for the FontAtlas and most plain image
        // sampling.
        SamplerSlot RegisterLinearClampSampler();
        // Shapes a UTF-8 string into glyph quads at (x, y) baseline. Calls
        // FontAtlas::Ensure for each codepoint (rasterising on first use),
        // emits one GlyphItem per visible glyph, returns the count written.
        // Use this to fill a GlyphItem buffer that you then dispatch.
        // Cursor advances along +X. No line-wrap, no kerning — single line.
        std::uint32_t ShapeText(Font& font, float pxSize,
                                float x, float baselineY,
                                std::string_view utf8,
@ -201,53 +245,79 @@ export namespace Crafter {
                                GlyphItem* out, std::uint32_t outCapacity,
                                float* outAdvance = nullptr);
        // Read after Initialize: the slot the font atlas was registered into.
        // 0xFFFF means "no atlas" (set fontAtlas before Initialize).
        std::uint16_t FontAtlasImageSlot()   const noexcept { return fontAtlasImageSlot_; }
        std::uint16_t FontAtlasSamplerSlot() const noexcept { return fontAtlasSamplerSlot_; }
        // Heap slot whose descriptor in each per-frame heap points at that
        // frame's swapchain image. Other passes (e.g. a ray-tracing pass
        // that wants to render the world directly into the swapchain) can
        // write to the same image by referencing this slot. Order in
        // window.passes controls compositing — push such passes BEFORE
        // the UI pass so UI overlays render on top.
        std::uint16_t OutImageSlot() const noexcept { return outImageSlot_; }
    private:
        Window* window_ = nullptr;
-        GraphicsDescriptorHeap* heap_ = nullptr;
+        DescriptorHeapVulkan* heap_ = nullptr;
        // One image slot used for the swapchain output. In each per-frame
        // heap, that slot points at THAT frame's swapchain image. So the
        // shader's `uiImages[hdr.outImage]` is always the current frame's
        // swapchain image regardless of which heap is bound.
        ImageSlot outImageSlot_;
        // Stable VkImageViewCreateInfos for the descriptor heap to ingest.
        // These must outlive the write call.
        VkImageViewCreateInfo atlasViewCreateInfo_{};
        ImageSlot   fontAtlasImageSlot_;
        SamplerSlot fontAtlasSamplerSlot_;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
        VkImageViewCreateInfo atlasViewCreateInfo_{};
        bool firstDispatchThisFrame_ = true;
 #endif
        // Subscription to window.onResize. Each resize destroys the old
        // swapchain images, so the per-frame heap entries we wrote at
        // outImageSlot_ now reference dangling VkImage handles. The
        // listener re-writes them and flushes the descriptor heaps.
        Crafter::EventListener<void> resizeSub_;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
        void WriteSwapchainDescriptors();
        void WriteFontAtlasDescriptor();
        // Helper used by RegisterBuffer template (defined in impl). Writes the
        // address-range descriptor at `slot` into all per-frame heaps.
        void WriteBufferDescriptor(std::uint16_t slot, VkDeviceAddress address, std::uint32_t size);
        // Helper used by RegisterImage template — writes a sampled-image at
        // `slot` referring to a stable VkImageViewCreateInfo (caller stores).
        void WriteSampledImageDescriptor(std::uint16_t slot,
                                         const VkImageViewCreateInfo& viewInfo,
                                         VkImageLayout layout);
 #endif
    };
    // ─── template-method implementations ────────────────────────────────
    template<typename T, bool Mapped>
-    BufferSlot UIRenderer::RegisterBuffer(GraphicsBuffer<T, Mapped>& buffer) {
+    BufferSlot UIRenderer::RegisterBuffer(VulkanBuffer<T, Mapped>& buffer) {
        auto range = heap_->AllocateBufferSlots(1);
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
        WriteBufferDescriptor(range.firstElement, buffer.address, buffer.size);
-#else
+        // BufferSlot's operator uint16_t() folds in heap_->bufferStartElement,
-        heap_->bufferTable[range.firstElement] = buffer.handle;
+        // so callers receive the absolute heap index when they convert.
 #endif
        return BufferSlot{heap_, range.firstElement};
    }
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
    template<typename Pixel>
    ImageSlot UIRenderer::RegisterImage(ImageVulkan<Pixel>& image, VkFormat format,
                                        VkImageLayout layout) {
        auto range = heap_->AllocateImageSlots(1);
        // Build a stable view-create-info that lives as long as the heap reads
        // it. We co-locate it on the renderer for the font atlas; for arbitrary
        // user images we lean on the fact that vkWriteResourceDescriptorsEXT
        // copies the view descriptor immediately. (Validated by the heap spec:
        // the descriptor is materialised at write time, the create-info need
        // not persist past the call.)
        VkImageViewCreateInfo info {
            .sType    = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO,
            .image    = image.image,
@ -268,5 +338,5 @@ export namespace Crafter {
        WriteSampledImageDescriptor(range.firstElement, info, layout);
        return ImageSlot{heap_, range.firstElement};
    }
 #endif
 }
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-UIComponents.cppm
+++ b/interfaces/Crafter.Graphics-UIComponents.cppm
@ -16,7 +16,12 @@ You should have received a copy of the GNU Lesser General Public
 License along with this library; if not, write to the Free Software
 Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301 USA
 */
 module;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
 export module Crafter.Graphics:UIComponents;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 import :UI;
 import :Font;
@ -138,3 +143,4 @@ export namespace Crafter {
                   std::array<float, 4> tint = {1, 1, 1, 1},
                   std::array<float, 4> uv   = {0, 0, 1, 1});
 }
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-WebGPU.cppm
+++ b/interfaces/Crafter.Graphics-WebGPU.cppm
@ -1,85 +0,0 @@
 /*
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 // JS bridge declarations for the DOM-mode WebGPU backend. Each function
 // corresponds to one entry in `additional/dom-webgpu.js`. Handles are
 // opaque uint32 cookies into the JS-side handle tables.
 export module Crafter.Graphics:WebGPU;
 #ifdef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 export namespace Crafter {
    using WebGPUBufferRef         = std::uint32_t;
    using WebGPUTextureRef        = std::uint32_t;
    using WebGPUSamplerRef        = std::uint32_t;
    using WebGPUCommandEncoderRef = std::uint32_t; // unused as a real handle; just a marker type for portability
 }
 namespace Crafter::WebGPU {
    __attribute__((import_module("env"), import_name("wgpuGetCanvasWidth")))
    extern "C" std::int32_t wgpuGetCanvasWidth();
    __attribute__((import_module("env"), import_name("wgpuGetCanvasHeight")))
    extern "C" std::int32_t wgpuGetCanvasHeight();
    __attribute__((import_module("env"), import_name("wgpuSurfaceWidth")))
    extern "C" std::int32_t wgpuSurfaceWidth();
    __attribute__((import_module("env"), import_name("wgpuSurfaceHeight")))
    extern "C" std::int32_t wgpuSurfaceHeight();
    __attribute__((import_module("env"), import_name("wgpuInit")))
    extern "C" void wgpuInit();
    __attribute__((import_module("env"), import_name("wgpuCreateBuffer")))
    extern "C" std::uint32_t wgpuCreateBuffer(std::int32_t byteSize);
    __attribute__((import_module("env"), import_name("wgpuWriteBuffer")))
    extern "C" void wgpuWriteBuffer(std::uint32_t handle, const void* srcPtr, std::int32_t byteSize);
    __attribute__((import_module("env"), import_name("wgpuDestroyBuffer")))
    extern "C" void wgpuDestroyBuffer(std::uint32_t handle);
    __attribute__((import_module("env"), import_name("wgpuCreateAtlasTexture")))
    extern "C" std::uint32_t wgpuCreateAtlasTexture(std::int32_t w, std::int32_t h);
    __attribute__((import_module("env"), import_name("wgpuWriteAtlasRegion")))
    extern "C" void wgpuWriteAtlasRegion(std::uint32_t handle, const void* srcPtr,
                                         std::int32_t srcW, std::int32_t srcH,
                                         std::int32_t srcBytesPerRow,
                                         std::int32_t dstX, std::int32_t dstY,
                                         std::int32_t copyW, std::int32_t copyH);
    __attribute__((import_module("env"), import_name("wgpuDestroyTexture")))
    extern "C" void wgpuDestroyTexture(std::uint32_t handle);
    __attribute__((import_module("env"), import_name("wgpuCreateLinearClampSampler")))
    extern "C" std::uint32_t wgpuCreateLinearClampSampler();
    __attribute__((import_module("env"), import_name("wgpuFrameBegin")))
    extern "C" void wgpuFrameBegin();
    __attribute__((import_module("env"), import_name("wgpuFrameEnd")))
    extern "C" void wgpuFrameEnd();
    __attribute__((import_module("env"), import_name("wgpuDispatchQuads")))
    extern "C" void wgpuDispatchQuads(std::uint32_t itemsHandle, const void* headerPtr,
                                      std::int32_t gx, std::int32_t gy);
    __attribute__((import_module("env"), import_name("wgpuDispatchCircles")))
    extern "C" void wgpuDispatchCircles(std::uint32_t itemsHandle, const void* headerPtr,
                                        std::int32_t gx, std::int32_t gy);
    __attribute__((import_module("env"), import_name("wgpuDispatchImages")))
    extern "C" void wgpuDispatchImages(std::uint32_t itemsHandle, const void* headerPtr,
                                       std::int32_t gx, std::int32_t gy,
                                       std::uint32_t texHandle, std::uint32_t sampHandle);
    __attribute__((import_module("env"), import_name("wgpuDispatchText")))
    extern "C" void wgpuDispatchText(std::uint32_t itemsHandle, const void* headerPtr,
                                     std::int32_t gx, std::int32_t gy,
                                     std::uint32_t atlasHandle, std::uint32_t sampHandle);
    // ─── custom user-authored compute shaders ───────────────────────────
    __attribute__((import_module("env"), import_name("wgpuLoadCustomShader")))
    extern "C" std::uint32_t wgpuLoadCustomShader(const void* wgslPtr, std::int32_t wgslLen,
                                                  const void* bindingsPtr, std::int32_t bindingsCount);
    __attribute__((import_module("env"), import_name("wgpuDispatchCustom")))
    extern "C" void wgpuDispatchCustom(std::uint32_t pipelineHandle,
                                       const void* pushPtr, std::int32_t pushBytes,
                                       const void* handlesPtr, std::int32_t handlesCount,
                                       std::int32_t gx, std::int32_t gy, std::int32_t gz);
 }
 #endif // CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-WebGPUBuffer.cppm
+++ b/interfaces/Crafter.Graphics-WebGPUBuffer.cppm
@ -1,85 +0,0 @@
 /*
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 // WebGPU buffer wrapper — DOM-mode parallel to VulkanBuffer<T, Mapped>.
 // Holds a JS-side GPUBuffer handle + (when Mapped) a wasm-memory staging
 // array. `.value` points to the staging memory; the user writes into it
 // directly, and `.Flush(cmd)` copies to the GPU via queue.writeBuffer.
 export module Crafter.Graphics:WebGPUBuffer;
 #ifdef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 import :WebGPU;
 export namespace Crafter {
    class WebGPUBufferBase {
    public:
        WebGPUBufferRef handle = 0;
        std::uint32_t size = 0;   // bytes
    };
    template<typename T>
    class WebGPUBufferMapped {
    public:
        T* value = nullptr;
    };
    class WebGPUBufferMappedEmpty {};
    template<typename T, bool Mapped>
    using WebGPUBufferMappedConditional =
        std::conditional_t<Mapped, WebGPUBufferMapped<T>, WebGPUBufferMappedEmpty>;
    template<typename T, bool Mapped>
    class WebGPUBuffer : public WebGPUBufferBase, public WebGPUBufferMappedConditional<T, Mapped> {
    public:
        WebGPUBuffer() = default;
        WebGPUBuffer(const WebGPUBuffer&) = delete;
        WebGPUBuffer& operator=(const WebGPUBuffer&) = delete;
        WebGPUBuffer(WebGPUBuffer&& other) noexcept {
            handle = other.handle;
            size   = other.size;
            other.handle = 0;
            if constexpr (Mapped) {
                this->value = other.value;
                other.value = nullptr;
            }
        }
        void Create(std::uint32_t count) {
            size = static_cast<std::uint32_t>(count * sizeof(T));
            handle = WebGPU::wgpuCreateBuffer(static_cast<std::int32_t>(size));
            if constexpr (Mapped) {
                this->value = new T[count]();
            }
        }
        void Clear() {
            if (handle != 0) {
                WebGPU::wgpuDestroyBuffer(handle);
                handle = 0;
            }
            if constexpr (Mapped) {
                if (this->value) { delete[] this->value; this->value = nullptr; }
            }
        }
        void Resize(std::uint32_t count) {
            if (handle != 0) Clear();
            Create(count);
        }
        void Flush(WebGPUCommandEncoderRef /*cmd*/) requires(Mapped) {
            WebGPU::wgpuWriteBuffer(handle, this->value, static_cast<std::int32_t>(size));
        }
        void FlushDevice() requires(Mapped) {
            WebGPU::wgpuWriteBuffer(handle, this->value, static_cast<std::int32_t>(size));
        }
        ~WebGPUBuffer() { Clear(); }
    };
 }
 #endif // CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm
+++ b/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm
@ -1,77 +0,0 @@
 /*
 Crafter®.Graphics
 Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 // User-authored compute shader for DOM mode.
 //
 // Contract:
 //   - WGSL source authored by the user.
 //   - Group 0 binding 0 is reserved for the UIDispatchHeader uniform
 //     (with dynamic offset). The library writes it from the first 48
 //     bytes of the push data each dispatch.
 //   - Group 1 is reserved for the ping-pong textures: binding 0 is the
 //     storage `out` (texture_storage_2d<rgba8unorm, write>), binding 1
 //     is the sampled `prev` (texture_2d<f32>). The library auto-binds
 //     the right textures depending on the current ping-pong state.
 //   - Groups 2+ are user-defined. The user declares each binding via a
 //     UICustomBinding descriptor at Load time, naming:
 //        - the @group(N) and @binding(N) numbers,
 //        - the resource KIND (buffer / sampled texture / sampler),
 //        - the BYTE OFFSET in the per-dispatch push data where a
 //          uint32 heap slot index lives.
 //     At Dispatch time the renderer reads each declared slot out of
 //     push data, looks the GPU handle up in the heap (bufferTable /
 //     imageTable / samplerTable), and assembles the bind group.
 export module Crafter.Graphics:WebGPUComputeShader;
 #ifdef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 import :WebGPU;
 export namespace Crafter {
    enum class UICustomBindingKind : std::uint8_t {
        Buffer         = 0,   // read-only-storage SSBO, handle is a slot into heap.bufferTable
        SampledTexture = 1,   // sampled texture_2d<f32>, handle is a slot into heap.imageTable
        Sampler        = 2,   // filtering sampler, handle is a slot into heap.samplerTable
    };
    struct UICustomBinding {
        std::uint8_t  group;        // @group(N), must be >= 2 (0 and 1 are reserved)
        std::uint8_t  binding;      // @binding(N)
        UICustomBindingKind kind;
        std::uint8_t  _pad;
        std::uint32_t pushOffset;   // offset in push data where the slot uint32 lives
    };
    static_assert(sizeof(UICustomBinding) == 8);
    class WebGPUComputeShader {
    public:
        std::uint32_t pipelineHandle = 0;
        std::vector<UICustomBinding> customBindings;
        WebGPUComputeShader() = default;
        WebGPUComputeShader(const WebGPUComputeShader&) = delete;
        WebGPUComputeShader& operator=(const WebGPUComputeShader&) = delete;
        WebGPUComputeShader(WebGPUComputeShader&& o) noexcept
            : pipelineHandle(o.pipelineHandle),
              customBindings(std::move(o.customBindings)) {
            o.pipelineHandle = 0;
        }
        // Compile + link a custom compute shader. `wgsl` is the source
        // string; the library does NOT add anything to it — the user's
        // shader must declare @group(0)/@group(1) bindings matching the
        // contract above. `bindings` lists every additional resource
        // (groups 2+) that the renderer should bind at dispatch time.
        void Load(std::string_view wgsl,
                  std::span<const UICustomBinding> bindings = {});
        // Path-based overload for symmetry with the Vulkan ComputeShader.
        // Reads the file from disk (browser VFS) and forwards to Load(wgsl).
        void Load(const std::filesystem::path& wgslPath,
                  std::span<const UICustomBinding> bindings = {});
    };
 }
 #endif // CRAFTER_GRAPHICS_WINDOW_DOM
--- a/interfaces/Crafter.Graphics-Window.cppm
+++ b/interfaces/Crafter.Graphics-Window.cppm
@ -50,10 +50,6 @@ import std;
 import :Types;
 import :Keys;
 import Crafter.Event;
 #ifdef CRAFTER_GRAPHICS_WINDOW_DOM
 import :WebGPU;
 import :DescriptorHeapWebGPU;
 #endif
 export namespace Crafter {
    #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
@ -65,9 +61,6 @@ export namespace Crafter {
 	};
    struct RenderPass;
    struct DescriptorHeapVulkan;
    #else
    struct RenderPass;
    struct DescriptorHeapWebGPU;
    #endif
    struct Window {
@ -249,21 +242,12 @@ export namespace Crafter {
        DescriptorHeapVulkan* descriptorHeap = nullptr;
        std::optional<std::array<float, 4>> clearColor;
        #else
-        // DOM mode: the page IS the window. WebGPU device and canvas are
+        // DOM mode: the page IS the window. `numFrames` stays as a public
-        // owned JS-side (see additional/dom-webgpu.js); this struct just
+        // constant so cross-platform code can refer to Window::numFrames
-        // holds the per-Window state Crafter::Window users expect:
+        // without #ifdef'ing the reference; nothing else lives here yet.
-        // a list of render passes and a pointer to the descriptor heap.
+        // V2 (WebGPU compute) will hang its GPUContext / swapchain texture
        // members off this branch.
        static constexpr std::uint8_t numFrames = 1;
        std::uint32_t currentBuffer = 0;
        std::vector<RenderPass*> passes;
        DescriptorHeapWebGPU* descriptorHeap = nullptr;
        std::optional<std::array<float, 4>> clearColor;
        // DOM-mode StartInit/FinishInit are no-ops returning an opaque
        // command-buffer marker so cross-platform user code (HelloUI's
        // `auto init = window.StartInit();`) compiles unchanged.
        WebGPUCommandEncoderRef StartInit();
        void FinishInit();
        #endif
    };
 }
--- a/interfaces/Crafter.Graphics.cppm
+++ b/interfaces/Crafter.Graphics.cppm
@ -34,7 +34,6 @@ export import :Input;
 export import :Device;
 export import :Animation;
 export import :ForwardDeclarations;
 export import :GraphicsTypes;
 export import :Clipboard;
 // Vulkan-backed partitions — empty under DOM.
@ -62,7 +61,3 @@ export import :Decompress;
 export import :Dom;
 export import :DomEvents;
 export import :Router;
 export import :WebGPU;
 export import :WebGPUBuffer;
 export import :DescriptorHeapWebGPU;
 export import :WebGPUComputeShader;
--- a/project.cpp
+++ b/project.cpp
@ -131,14 +131,13 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
    // when its body is gated out. Vulkan-typed partitions stub to empty
    // modules under CRAFTER_GRAPHICS_WINDOW_DOM; the Dom/DomEvents/Router
    // partitions stub to empty modules in the opposite direction.
-    std::array<fs::path, 37> ifaces = {
+    std::array<fs::path, 32> ifaces = {
        "interfaces/Crafter.Graphics",
        "interfaces/Crafter.Graphics-Animation",
        "interfaces/Crafter.Graphics-Clipboard",
        "interfaces/Crafter.Graphics-ComputeShader",
        "interfaces/Crafter.Graphics-Decompress",
        "interfaces/Crafter.Graphics-DescriptorHeapVulkan",
        "interfaces/Crafter.Graphics-DescriptorHeapWebGPU",
        "interfaces/Crafter.Graphics-Device",
        "interfaces/Crafter.Graphics-Dom",
        "interfaces/Crafter.Graphics-DomEvents",
@ -146,7 +145,6 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
        "interfaces/Crafter.Graphics-FontAtlas",
        "interfaces/Crafter.Graphics-ForwardDeclarations",
        "interfaces/Crafter.Graphics-Gamepad",
        "interfaces/Crafter.Graphics-GraphicsTypes",
        "interfaces/Crafter.Graphics-ImageVulkan",
        "interfaces/Crafter.Graphics-Input",
        "interfaces/Crafter.Graphics-InputField",
@ -165,37 +163,29 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
        "interfaces/Crafter.Graphics-UIComponents",
        "interfaces/Crafter.Graphics-VulkanBuffer",
        "interfaces/Crafter.Graphics-VulkanTransition",
        "interfaces/Crafter.Graphics-WebGPU",
        "interfaces/Crafter.Graphics-WebGPUBuffer",
        "interfaces/Crafter.Graphics-WebGPUComputeShader",
        "interfaces/Crafter.Graphics-Window",
    };
    if (dom) {
-        // DOM impl set. UI-Shared.cpp is backend-agnostic; UI-WebGPU.cpp
+        // DOM impl set: only the files whose bodies do meaningful work
-        // is the DOM-only implementation of UIRenderer's GPU-touching
+        // under CRAFTER_GRAPHICS_WINDOW_DOM. Vulkan-only impls (Mesh,
-        // methods. Font / FontAtlas / UIComponents are now portable.
+        // ComputeShader, Font, etc.) stay out — their interface stubs
-        std::array<fs::path, 12> domImpls = {
+        // mean no link-side references; including the impl would just
        // be dead code with no Vulkan headers to compile against.
        std::array<fs::path, 6> domImpls = {
            "implementations/Crafter.Graphics-Clipboard",
            "implementations/Crafter.Graphics-Dom",
            "implementations/Crafter.Graphics-Font",
            "implementations/Crafter.Graphics-FontAtlas",
            "implementations/Crafter.Graphics-Gamepad",
            "implementations/Crafter.Graphics-Input",
            "implementations/Crafter.Graphics-Router",
            "implementations/Crafter.Graphics-UI-Shared",
            "implementations/Crafter.Graphics-UI-WebGPU",
            "implementations/Crafter.Graphics-UIComponents",
            "implementations/Crafter.Graphics-WebGPUComputeShader",
            "implementations/Crafter.Graphics-Window",
        };
        cfg.GetInterfacesAndImplementations(ifaces, domImpls);
        // JS glue shipped alongside the .wasm so the loader has the
        // env-import surface the Window/Dom bindings expect.
        cfg.files.emplace_back(fs::path("additional/dom-env.js"));
        cfg.files.emplace_back(fs::path("additional/dom-webgpu.js"));
    } else {
-        std::array<fs::path, 14> impls = {
+        std::array<fs::path, 13> impls = {
            "implementations/Crafter.Graphics-Clipboard",
            "implementations/Crafter.Graphics-ComputeShader",
            "implementations/Crafter.Graphics-Device",
@ -207,7 +197,6 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
            "implementations/Crafter.Graphics-Mesh",
            "implementations/Crafter.Graphics-RenderingElement3D",
            "implementations/Crafter.Graphics-UI",
            "implementations/Crafter.Graphics-UI-Shared",
            "implementations/Crafter.Graphics-UIComponents",
            "implementations/Crafter.Graphics-Window",
        };
		`@ -0,0 +1,2 @@`
							`#!/usr/bin/env sh`
							`caddy file-server --listen :8080 --root bin/HelloDom-wasm32-wasip1-native-native-df37fe0fe124fe57`