diff --git a/additional/dom-webgpu.js b/additional/dom-webgpu.js index 81eafe3..8706e90 100644 --- a/additional/dom-webgpu.js +++ b/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 -// has its own bgl1 instance even if the layout entries are identical). -function getGroup1BG(pipe) { - const key = `${pipe.bgl1.label || ""}/${state.outIsPing ? 1 : 0}/${state.width}x${state.height}`; +// 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: 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 out +// group 1 binding 1 — texture_2d 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) => { diff --git a/examples/CustomShader/inverse-circle.comp.wgsl b/examples/CustomShader/inverse-circle.comp.wgsl new file mode 100644 index 0000000..81341ee --- /dev/null +++ b/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 out (auto) +// group 1 binding 1 — texture_2d 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 hdr : UIDispatchHeader; +@group(1) @binding(0) var outTex : texture_storage_2d; +@group(1) @binding(1) var prevTex : texture_2d; + +struct InverseCircleItem { + centerRadius: vec4, +}; +@group(2) @binding(0) var items : array; + +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) gid: vec3) { + if (gid.x >= hdr.surfaceW || gid.y >= hdr.surfaceH) { return; } + let coord = vec2(i32(gid.x), i32(gid.y)); + let sp = vec2(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(mix(dst.rgb, vec3(1.0) - dst.rgb, coverage), dst.a); + } + textureStore(outTex, coord, dst); +} diff --git a/examples/CustomShader/main.cpp b/examples/CustomShader/main.cpp index 9469b25..4d798c7 100644 --- a/examples/CustomShader/main.cpp +++ b/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 // 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 - // the library — there is no privileged path. - ComputeShader inverseCircle; + // Load the user-authored shader. On native it's an offline-compiled + // SPIR-V from .comp.glsl; on DOM it's WGSL source compiled at + // 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(offsetof(Crafter::UIDispatchHeader, itemBuffer)) }, + }; + inverseCircle.Load(std::filesystem::path("inverse-circle.comp.wgsl"), invBindings); +#endif // User-owned buffers. - VulkanBuffer quadsBuf; - VulkanBuffer invBuf; + GraphicsBuffer quadsBuf; + GraphicsBuffer 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 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; diff --git a/examples/CustomShader/project.cpp b/examples/CustomShader/project.cpp index 9e611a8..f9d2f29 100644 --- a/examples/CustomShader/project.cpp +++ b/examples/CustomShader/project.cpp @@ -4,6 +4,14 @@ namespace fs = std::filesystem; using namespace Crafter; extern "C" Configuration CrafterBuildProject(std::span 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(args.begin(), args.end()), @@ -13,6 +21,14 @@ extern "C" Configuration CrafterBuildProject(std::span 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 a std::array 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; } diff --git a/examples/HelloUI/project.cpp b/examples/HelloUI/project.cpp index dc03f3c..409d66d 100644 --- a/examples/HelloUI/project.cpp +++ b/examples/HelloUI/project.cpp @@ -25,6 +25,10 @@ extern "C" Configuration CrafterBuildProject(std::span 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 }; diff --git a/implementations/Crafter.Graphics-UI-WebGPU.cpp b/implementations/Crafter.Graphics-UI-WebGPU.cpp index 4050974..6e619c5 100644 --- a/implementations/Crafter.Graphics-UI-WebGPU.cpp +++ b/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 handles; + handles.reserve(shader.customBindings.size()); + const std::uint8_t* p = static_cast(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(pushBytes), + handles.data(), static_cast(handles.size()), + static_cast(gx), + static_cast(gy), + static_cast(gz)); +} diff --git a/implementations/Crafter.Graphics-UI.cpp b/implementations/Crafter.Graphics-UI.cpp index c7d255c..e8ce871 100644 --- a/implementations/Crafter.Graphics-UI.cpp +++ b/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_) { diff --git a/implementations/Crafter.Graphics-WebGPUComputeShader.cpp b/implementations/Crafter.Graphics-WebGPUComputeShader.cpp new file mode 100644 index 0000000..5a8e554 --- /dev/null +++ b/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 bindings) { + customBindings.assign(bindings.begin(), bindings.end()); + pipelineHandle = WebGPU::wgpuLoadCustomShader( + wgsl.data(), + static_cast(wgsl.size()), + customBindings.data(), + static_cast(customBindings.size()) + ); +} + +void WebGPUComputeShader::Load(const std::filesystem::path& wgslPath, + std::span 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(size), '\0'); + f.read(src.data(), size); + Load(std::string_view{src}, bindings); +} diff --git a/interfaces/Crafter.Graphics-UI.cppm b/interfaces/Crafter.Graphics-UI.cppm index 00bf381..6dc0d36 100644 --- a/interfaces/Crafter.Graphics-UI.cppm +++ b/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 clipRectPx = {0.0f, 0.0f, 1e9f, 1e9f}); -#ifndef CRAFTER_GRAPHICS_WINDOW_DOM - // Generic dispatch — user-authored shaders. Vulkan-only in v1; on DOM - // the WebGPU side has no bindless and would need per-shader bind-group - // declaration. See plan section 3b for the design path. - void Dispatch(GraphicsCommandBuffer cmd, const ComputeShader& shader, + // Generic dispatch for user-authored shaders. On Vulkan, `shader` is + // a SPIR-V compute pipeline (bindless via VK_EXT_descriptor_heap, so + // any resource indices baked into push data resolve through the + // 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); -#endif // Allocates a heap slot for the buffer and registers the GPU handle. // Returns a move-only BufferSlot RAII handle. diff --git a/interfaces/Crafter.Graphics-WebGPU.cppm b/interfaces/Crafter.Graphics-WebGPU.cppm index d6d3a0b..6831322 100644 --- a/interfaces/Crafter.Graphics-WebGPU.cppm +++ b/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 diff --git a/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm b/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm index 62eda86..fac065f 100644 --- a/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm +++ b/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 -// 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 -// declare `WebGPUComputeShader drawQuads;` members for symmetry with the -// Vulkan side, but `Load()` and `Dispatch()` are intentionally absent — -// the DispatchQuads / DispatchCircles / etc convenience methods on -// UIRenderer route directly to the JS bridge. +// 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), binding 1 +// is the sampled `prev` (texture_2d). 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 { - // Marker only; pipelines live JS-side per dispatchStandard in - // dom-webgpu.js. No state required. + enum class UICustomBindingKind : std::uint8_t { + Buffer = 0, // read-only-storage SSBO, handle is a slot into heap.bufferTable + SampledTexture = 1, // sampled texture_2d, 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 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 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 bindings = {}); }; } #endif // CRAFTER_GRAPHICS_WINDOW_DOM diff --git a/project.cpp b/project.cpp index d182ef3..06f64ac 100644 --- a/project.cpp +++ b/project.cpp @@ -175,7 +175,7 @@ extern "C" Configuration CrafterBuildProject(std::span 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 domImpls = { + std::array domImpls = { "implementations/Crafter.Graphics-Clipboard", "implementations/Crafter.Graphics-Dom", "implementations/Crafter.Graphics-Font", @@ -186,6 +186,7 @@ extern "C" Configuration CrafterBuildProject(std::span 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);