custom shader webgpu

additional/dom-webgpu.js

@@ -413,16 +413,18 @@ const hdrBG = {
 };
 
 // Group 1 changes between dispatches because `out` and `prev` swap on the
-// ping-pong. Cache by current out-is-ping bool + pipeline (each pipeline
+// ping-pong. Cached by current ping-pong direction and texture size; the
-// has its own bgl1 instance even if the layout entries are identical).
+// stored bind group is reusable across all pipelines that share a
-function getGroup1BG(pipe) {
+// layout-compatible bgl1 (all standard pipelines and custom shaders do,
-    const key = `${pipe.bgl1.label || ""}/${state.outIsPing ? 1 : 0}/${state.width}x${state.height}`;
+// 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: pipe.bgl1,
+        layout: bgl1,
         entries: [
             { binding: 0, resource: outView },
             { binding: 1, resource: prevView },
@@ -662,7 +664,7 @@ function dispatchStandard(pipe, hdrBindGroup, headerPtr, gx, gy, itemsHandle, gr
     const off = writeHeader(headerPtr);
     state.pass.setPipeline(pipe.pipeline);
     state.pass.setBindGroup(0, hdrBindGroup, [off]);
-    state.pass.setBindGroup(1, getGroup1BG(pipe));
+    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);
@@ -686,6 +688,136 @@ env.wgpuDispatchText = (itemsHandle, headerPtr, gx, gy, 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;
@@ -708,6 +840,8 @@ window.crafter_wgpu_debug = () => ({
     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) => {

examples/CustomShader/inverse-circle.comp.wgsl

@@ -0,0 +1,58 @@
+// 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);
+}

examples/CustomShader/main.cpp

@@ -2,8 +2,17 @@
 // 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;
@@ -11,18 +20,22 @@ import std;
 using namespace Crafter;
 
 // Application-side item POD. Matches `struct InverseCircleItem { vec4
-// centerRadius; }` in inverse-circle.comp.glsl byte-for-byte.
+// centerRadius; }` in inverse-circle.comp.{glsl,wgsl} 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
 
-    VkCommandBuffer init = window.StartInit();
+    auto init = window.StartInit();
 
-    DescriptorHeapVulkan heap;
+    GraphicsDescriptorHeap heap;
     heap.Initialize(/*images*/ 8, /*buffers*/ 8, /*samplers*/ 4);
     window.descriptorHeap = &heap;
 
@@ -30,26 +43,46 @@ int main() {
     ui.Initialize(window, heap, init);
     window.passes.push_back(&ui);
 
-    // Load the user-authored shader. Same wrapper as the four shipped with
+    // Load the user-authored shader. On native it's an offline-compiled
-    // the library — there is no privileged path.
+    // SPIR-V from .comp.glsl; on DOM it's WGSL source compiled at
-    ComputeShader inverseCircle;
+    // startup by the device. The DOM Load takes an extra `bindings` arg
+    // 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.
-    VulkanBuffer<QuadItem,           true> quadsBuf;
+    GraphicsBuffer<QuadItem,          true> quadsBuf;
-    VulkanBuffer<InverseCircleItem,  true> invBuf;
+    GraphicsBuffer<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) {
-        VkCommandBuffer cmd = a.cmd;
+        auto cmd = a.cmd;
 
         Rect canvas = Rect::FromWindow(window);
 
@@ -82,15 +115,24 @@ 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.
+        // SHADER_WRITE → SHADER_READ|WRITE barrier automatically on
+        // 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;

examples/CustomShader/project.cpp

@@ -4,6 +4,14 @@ 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()),
@@ -13,6 +21,14 @@ 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 };
 
@@ -20,6 +36,14 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
     std::array<fs::path, 1> impls = { "main" };
     cfg.GetInterfacesAndImplementations(ifaces, impls);
 
-    cfg.shaders.emplace_back(fs::path("inverse-circle.comp.glsl"), std::string("main"), ShaderType::Compute);
+    if (isWasm) {
+        // 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;
 }

examples/HelloUI/project.cpp

@@ -25,6 +25,10 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
     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 };

implementations/Crafter.Graphics-UI-WebGPU.cpp

@@ -138,3 +138,43 @@ SamplerSlot UIRenderer::RegisterLinearClampSampler() {
     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));
+}

implementations/Crafter.Graphics-UI.cpp

@@ -178,7 +178,7 @@ void UIRenderer::DispatchText(GraphicsCommandBuffer cmd, std::uint32_t bufferSlo
 
 // ─── generic Dispatch (with barrier) ────────────────────────────────────
 
-void UIRenderer::Dispatch(GraphicsCommandBuffer cmd, const ComputeShader& shader,
+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) {
     if (!firstDispatchThisFrame_) {

implementations/Crafter.Graphics-WebGPUComputeShader.cpp

@@ -0,0 +1,41 @@
+/*
+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);
+}

interfaces/Crafter.Graphics-UI.cppm

@@ -168,14 +168,17 @@ export namespace Crafter {
         void DispatchText(GraphicsCommandBuffer cmd, std::uint32_t bufferSlot, std::uint32_t itemCount,
                           std::array<float,4> clipRectPx = {0.0f, 0.0f, 1e9f, 1e9f});
 
-#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
+        // Generic dispatch for user-authored shaders. On Vulkan, `shader` is
-        // Generic dispatch — user-authored shaders. Vulkan-only in v1; on DOM
+        // a SPIR-V compute pipeline (bindless via VK_EXT_descriptor_heap, so
-        // the WebGPU side has no bindless and would need per-shader bind-group
+        // any resource indices baked into push data resolve through the
-        // declaration. See plan section 3b for the design path.
+        // global heap). On DOM, `shader` carries a UICustomBinding list
-        void Dispatch(GraphicsCommandBuffer cmd, const ComputeShader& shader,
+        // 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);
-#endif
 
         // Allocates a heap slot for the buffer and registers the GPU handle.
         // Returns a move-only BufferSlot RAII handle.

interfaces/Crafter.Graphics-WebGPU.cppm

@@ -71,5 +71,15 @@ namespace Crafter::WebGPU {
     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

interfaces/Crafter.Graphics-WebGPUComputeShader.cppm

@@ -4,22 +4,74 @@ Copyright (C) 2026 Catcrafts®
 catcrafts.net
 */
 
-// Placeholder ComputeShader for DOM mode. The four standard UI pipelines
+// User-authored compute shader for DOM mode.
-// are compiled JS-side at startup (see additional/dom-webgpu.js); the C++
+//
-// side never sees a WGSL handle. This type exists so UIRenderer can
+// Contract:
-// declare `WebGPUComputeShader drawQuads;` members for symmetry with the
+//   - WGSL source authored by the user.
-// Vulkan side, but `Load()` and `Dispatch()` are intentionally absent —
+//   - Group 0 binding 0 is reserved for the UIDispatchHeader uniform
-// the DispatchQuads / DispatchCircles / etc convenience methods on
+//     (with dynamic offset). The library writes it from the first 48
-// UIRenderer route directly to the JS bridge.
+//     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 {
-    struct WebGPUComputeShader {
+    enum class UICustomBindingKind : std::uint8_t {
-        // Marker only; pipelines live JS-side per dispatchStandard in
+        Buffer         = 0,   // read-only-storage SSBO, handle is a slot into heap.bufferTable
-        // dom-webgpu.js. No state required.
+        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

project.cpp

@@ -175,7 +175,7 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
         // DOM impl set. UI-Shared.cpp is backend-agnostic; UI-WebGPU.cpp
         // is the DOM-only implementation of UIRenderer's GPU-touching
         // methods. Font / FontAtlas / UIComponents are now portable.
-        std::array<fs::path, 11> domImpls = {
+        std::array<fs::path, 12> domImpls = {
             "implementations/Crafter.Graphics-Clipboard",
             "implementations/Crafter.Graphics-Dom",
             "implementations/Crafter.Graphics-Font",
@@ -186,6 +186,7 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
             "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);