test(webgpu): HDRBloom example exercising the HDR post-process primitives (#27)

RT→rgba16float→threshold→blur→composite, end-to-end proof of the three
primitives. threshold/blur run from onBeforeUpdate (one submit each) so the
storage-write→sampled-read dependency gets WebGPU's per-submit barrier
(there is none between dispatches within a compute pass); the composite is a
UI custom shader so it owns the canvas ping-pong. Documents the Vulkan
symmetry (gap 4): the native present path records passes generically and
barriers between them, so the same chain is wireable today.

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

WAVEFRONT-DESIGN.md

@@ -42,6 +42,16 @@ Accum buffer is linear. Optional user `WebGPURTStage::Resolve` entry
 `resolve_main(coord:vec2<u32>, hdr:vec4<f32>)->vec4<f32>`. None → passthrough.
 VulkanTriangle: no resolve (exact match). Sponza: resolve does Reinhard+gamma.
 
+### HDR output target (issue #27)
+`PipelineRTWebGPU::Init(..., hdrOutputFormat)` (default `RGBA8Unorm` = the
+canvas ping-pong path, unchanged) can instead point RESOLVE at a user
+`rgba16float` storage texture (`RTPass::outTexHandle`). The JS side swaps
+binding(6)'s WGSL declaration + bind-group-layout format and skips the
+ping-pong flip, since the canvas is untouched. With no resolve shader the
+default passthrough writes raw linear radiance — the HDR input an app's own
+bloom/composite chain (threshold → blur → tonemap → swapchain) reads. See
+`examples/HDRBloom`.
+
 ## Indirect dispatch (Phase 2 de-risk)
 Prove `dispatchWorkgroupsIndirect` + cross-pass atomic visibility with a toy
 "emit N → dispatch N" before wiring real kernels. WebGPU inserts an implicit

examples/HDRBloom/blur.comp.wgsl

@@ -0,0 +1,33 @@
+// HDRBloom blur pass (PlainComputeShader). Box-blurs the thresholded bloom
+// mip into a second rgba16float target — demonstrating an N-target float
+// chain: this pass SAMPLES the texture the threshold pass wrote (storage
+// write → sampled read across a submit barrier) and WRITES another
+// user-owned float storage texture. Two such targets prove the mip-pyramid
+// shape a real bloom needs.
+//
+// Layout mirrors threshold.comp.wgsl (src = bloom mip A, dst = bloom mip B).
+
+struct Dim { w: u32, h: u32, _0: u32, _1: u32 };
+@group(0) @binding(0) var<uniform> dim : Dim;
+@group(1) @binding(0) var src : texture_2d<f32>;
+@group(1) @binding(1) var dst : texture_storage_2d<rgba16float, write>;
+
+const R: i32 = 6;   // box radius → 13×13 kernel; a wide, cheap glow
+
+@compute @workgroup_size(8, 8, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    if (gid.x >= dim.w || gid.y >= dim.h) { return; }
+    let maxX = i32(dim.w) - 1;
+    let maxY = i32(dim.h) - 1;
+    var sum = vec3<f32>(0.0);
+    var count = 0.0;
+    for (var dy = -R; dy <= R; dy = dy + 1) {
+        for (var dx = -R; dx <= R; dx = dx + 1) {
+            let x = clamp(i32(gid.x) + dx, 0, maxX);
+            let y = clamp(i32(gid.y) + dy, 0, maxY);
+            sum = sum + textureLoad(src, vec2<i32>(x, y), 0).rgb;
+            count = count + 1.0;
+        }
+    }
+    textureStore(dst, vec2<i32>(i32(gid.x), i32(gid.y)), vec4<f32>(sum / count, 1.0));
+}

examples/HDRBloom/closesthit.wgsl

@@ -0,0 +1,45 @@
+// HDRBloom closest-hit (runs in SHADE). Half the cubes are "emitters" that
+// accumulate a strongly super-1.0 radiance (the HDR signal a bloom pass
+// extracts); the rest are dim Lambert-shaded fillers near/below 1.0 that the
+// threshold rejects. The linear accumulator is what RESOLVE writes into the
+// rgba16float scene target — no tonemap here.
+//
+// Payload declared here so the assembler sees it before wfPayload / SHADE.
+struct Payload {
+    color: vec3<f32>,
+};
+
+const SUN_DIR_TO_LIGHT: vec3<f32> = vec3<f32>(0.40, 0.85, 0.35);
+const AMBIENT_COLOR:    vec3<f32> = vec3<f32>(0.10, 0.11, 0.16);
+
+// Distinct per-instance hue so emitters read as different coloured glows.
+fn instanceAlbedo(i: u32) -> vec3<f32> {
+    let h = i * 2654435761u;
+    return vec3<f32>(
+        0.35 + 0.6 * f32((h >>  0u) & 255u) / 255.0,
+        0.35 + 0.6 * f32((h >>  8u) & 255u) / 255.0,
+        0.35 + 0.6 * f32((h >> 16u) & 255u) / 255.0);
+}
+
+fn closesthit_main(ray: RayDesc, hit: HitInfo, payload: ptr<function, Payload>) {
+    let meshRec = meshRecords[tlasEntries[hit.instanceId].blasMeshIdx];
+    let verts = _rtFetchTri(meshRec, hit.primitiveId);
+    let nObj = normalize(cross(verts[1] - verts[0], verts[2] - verts[0]));
+    let nWorld = normalize(vec3<f32>(
+        dot(hit.objectToWorldR0.xyz, nObj),
+        dot(hit.objectToWorldR1.xyz, nObj),
+        dot(hit.objectToWorldR2.xyz, nObj)));
+
+    let albedo  = instanceAlbedo(hit.customIndex);
+    let viewDir = -ray.direction;
+    let nFacing = select(-nWorld, nWorld, dot(nWorld, viewDir) > 0.0);
+    let nDotL   = max(0.0, dot(nFacing, normalize(SUN_DIR_TO_LIGHT)));
+
+    // Every third cube is a bright emitter (HDR, > 1.0); the rest stay dim.
+    if ((hit.customIndex % 3u) == 0u) {
+        // View-independent emission so the whole face glows uniformly.
+        rtAccumulate(albedo * 9.0);
+    } else {
+        rtAccumulate(albedo * (AMBIENT_COLOR + vec3<f32>(0.55 * nDotL)));
+    }
+}

examples/HDRBloom/composite.comp.wgsl

@@ -0,0 +1,53 @@
+// HDRBloom composite (UI custom shader, dispatched via UIRenderer). Runs
+// inside the per-frame UI compute pass, so it owns the ping-pong `out`
+// texture that gets blitted to the canvas — the app's composite→swapchain
+// step. It reads the linear HDR scene (group 2 binding 0) and the blurred
+// bloom mip (group 2 binding 1, sampled through a filtering sampler — float
+// textures are filterable), adds them, tonemaps (Reinhard) and gamma-
+// corrects, then writes the rgba8unorm canvas.
+//
+//   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, unused)
+//   group 2 binding 0 — texture_2d<f32>  hdr scene    (heap slot in push)
+//   group 2 binding 1 — texture_2d<f32>  bloom mip     (heap slot in push)
+//   group 2 binding 2 — sampler          linear clamp  (heap slot in push)
+
+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>;
+@group(2) @binding(0) var hdrTex   : texture_2d<f32>;
+@group(2) @binding(1) var bloomTex : texture_2d<f32>;
+@group(2) @binding(2) var samp     : sampler;
+
+const BLOOM_STRENGTH: f32 = 1.0;
+
+@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 scene = textureLoad(hdrTex, coord, 0).rgb;
+    let uv    = (vec2<f32>(f32(gid.x), f32(gid.y)) + vec2<f32>(0.5))
+              / vec2<f32>(f32(hdr.surfaceW), f32(hdr.surfaceH));
+    let bloom = textureSampleLevel(bloomTex, samp, uv, 0.0).rgb;
+
+    let lit    = scene + bloom * BLOOM_STRENGTH;
+    let mapped = lit / (lit + vec3<f32>(1.0));        // Reinhard tonemap
+    let g      = pow(mapped, vec3<f32>(1.0 / 2.2));   // gamma 2.2
+    textureStore(outTex, coord, vec4<f32>(g, 1.0));
+}

examples/HDRBloom/main.cpp

@@ -0,0 +1,244 @@
+// HDRBloom — cross-backend-shaped HDR post-process on the WebGPU/DOM
+// backend, exercising the primitives added for issue #27:
+//
+//   1. rgba16float runtime textures (StorageImage2D) with STORAGE + SAMPLED
+//      usage — the scene + bloom-mip targets.
+//   2. A write-only StorageTexture binding kind (UICustomBindingKind::
+//      StorageTexture) so compute passes write those float targets, plus
+//      float-texture sampling via SampledTexture.
+//   3. PipelineRTWebGPU's RESOLVE writing linear radiance into a user
+//      rgba16float texture (hdrOutputFormat) instead of the canvas.
+//
+// Pipeline shape (the same one a Vulkan bloom would use):
+//   RT  → linear rgba16float scene
+//   threshold (compute, sample float → write float)
+//   blur      (compute, sample float → write float)
+//   composite (UI custom shader: scene + bloom → tonemap+gamma → canvas)
+//
+// The threshold/blur passes run from onBeforeUpdate so each lands on its own
+// queue submit — WebGPU's per-submit ordering gives the storage-write →
+// sampled-read barrier the chain needs (there is no barrier between
+// dispatches within one compute pass). The composite runs in-frame as a UI
+// custom shader so it owns the canvas ping-pong. Bloom is therefore one
+// frame behind the sharp scene, which is imperceptible for this static view.
+//
+// WebGPU/DOM only — the wavefront tracer is the WebGPU software RT path.
+
+#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
+int main() { return 0; }   // native bloom is wireable today (see issue gap 4)
+#else
+
+#include <cstddef>  // offsetof
+
+import Crafter.Graphics;
+import Crafter.Math;
+import Crafter.Event;
+import std;
+
+using namespace Crafter;
+namespace fs = std::filesystem;
+
+namespace {
+    constexpr int   kGrid    = 3;
+    constexpr float kSpacing = 2.5f;
+    constexpr float kHalf    = 0.5f;
+
+    struct CameraGPU {
+        float origin[3];   float pad0;
+        float right[3];    float tanHalf;
+        float up[3];       float aspect;
+        float forward[3];  float pad1;
+    };
+    static_assert(sizeof(CameraGPU) == 64);
+
+    struct Dim { std::uint32_t w, h, _0, _1; };
+
+    // Composite push: standard header + the three heap slots the UI custom
+    // shader's group(2) bindings resolve through.
+    struct CompositePush {
+        UIDispatchHeader hdr;
+        std::uint32_t    hdrSlot;
+        std::uint32_t    bloomSlot;
+        std::uint32_t    sampSlot;
+        std::uint32_t    _pad;
+    };
+}
+
+int main() {
+    Device::Initialize();
+    static Window window(1280, 720, "HDRBloom");
+    auto cmd = window.StartInit();
+
+    DescriptorHeapWebGPU heap;
+    heap.Initialize(/*images*/ 4, /*buffers*/ 4, /*samplers*/ 2);
+    window.descriptorHeap = &heap;
+
+    const std::uint16_t W = static_cast<std::uint16_t>(window.width);
+    const std::uint16_t H = static_cast<std::uint16_t>(window.height);
+
+    // ── RT pipeline: HDR output (RESOLVE → rgba16float) ────────────────
+    std::array<WebGPUShader, 3> shaders {{
+        WebGPUShader(fs::path("raygen.wgsl"),     "raygen_main",     WebGPURTStage::Raygen),
+        WebGPUShader(fs::path("miss.wgsl"),       "miss_main",       WebGPURTStage::Miss),
+        WebGPUShader(fs::path("closesthit.wgsl"), "closesthit_main", WebGPURTStage::ClosestHit),
+    }};
+    ShaderBindingTableWebGPU sbt;
+    sbt.Init(shaders);
+
+    std::array<RTShaderGroup, 1> raygenGroups {{ { .type = RTShaderGroupType::General,           .generalShader = 0 } }};
+    std::array<RTShaderGroup, 1> missGroups   {{ { .type = RTShaderGroupType::General,           .generalShader = 1 } }};
+    std::array<RTShaderGroup, 1> hitGroups    {{ { .type = RTShaderGroupType::TrianglesHitGroup, .closestHitShader = 2 } }};
+
+    std::array<UICustomBinding, 1> rtBindings {{
+        { .group = 3, .binding = 0, .kind = UICustomBindingKind::Buffer, .pushOffset = 0 },
+    }};
+
+    PipelineRTWebGPU pipeline;
+    pipeline.Init(cmd, raygenGroups, missGroups, hitGroups, sbt, rtBindings,
+                  WebGPUTexelFormat::RGBA16Float);   // ← RESOLVE writes HDR
+
+    // ── Unit cube mesh. ────────────────────────────────────────────────
+    static std::array<Vector<float, 3, 3>, 8> verts {{
+        {-kHalf, -kHalf, -kHalf}, { kHalf, -kHalf, -kHalf},
+        { kHalf,  kHalf, -kHalf}, {-kHalf,  kHalf, -kHalf},
+        {-kHalf, -kHalf,  kHalf}, { kHalf, -kHalf,  kHalf},
+        { kHalf,  kHalf,  kHalf}, {-kHalf,  kHalf,  kHalf},
+    }};
+    static std::array<std::uint32_t, 36> indices {{
+        0,1,2, 0,2,3,   5,4,7, 5,7,6,   4,0,3, 4,3,7,
+        1,5,6, 1,6,2,   4,5,1, 4,1,0,   3,2,6, 3,6,7,
+    }};
+    static Mesh cube;
+    cube.Build(verts, indices, cmd);
+
+    WebGPUBuffer<CameraGPU, true> cameraBuf;
+    cameraBuf.Create(1);
+    static std::array<std::uint32_t, 1> rtHandles { cameraBuf.handle };
+
+    static std::vector<RenderingElement3D> renderers;
+    renderers.reserve(static_cast<std::size_t>(kGrid * kGrid * kGrid));
+    const float origin0 = -0.5f * static_cast<float>(kGrid - 1) * kSpacing;
+    for (int x = 0; x < kGrid; ++x)
+    for (int y = 0; y < kGrid; ++y)
+    for (int z = 0; z < kGrid; ++z) {
+        renderers.emplace_back();
+        RenderingElement3D& r = renderers.back();
+        auto& tx = r.instance.transform.matrix;
+        tx[0][0] = 1; tx[0][1] = 0; tx[0][2] = 0; tx[0][3] = origin0 + float(x) * kSpacing;
+        tx[1][0] = 0; tx[1][1] = 1; tx[1][2] = 0; tx[1][3] = origin0 + float(y) * kSpacing;
+        tx[2][0] = 0; tx[2][1] = 0; tx[2][2] = 1; tx[2][3] = origin0 + float(z) * kSpacing;
+        r.instance.instanceCustomIndex                    = static_cast<std::uint32_t>(renderers.size() - 1);
+        r.instance.mask                                   = 0xFF;
+        r.instance.instanceShaderBindingTableRecordOffset = 0;
+        r.instance.flags                                  = kRTGeometryInstanceForceOpaque;
+        r.instance.accelerationStructureReference         = cube.blasAddr;
+        RenderingElement3D::Add(&r);
+    }
+    RenderingElement3D::BuildTLAS(cmd, 0);
+
+    // ── HDR scene + bloom mip targets (rgba16float). ───────────────────
+    StorageImage2D hdrScene;  hdrScene.Create(W, H);
+    StorageImage2D bloomA;    bloomA.Create(W, H);
+    StorageImage2D bloomB;    bloomB.Create(W, H);
+
+    // Heap slots for the composite's sampled inputs + sampler.
+    ImageSlot   hdrSlot   = hdrScene.AllocateSlot(heap);
+    ImageSlot   bloomSlot = bloomB.AllocateSlot(heap);
+    SamplerSlot sampSlot  = AllocateLinearClampSampler(heap);
+
+    // ── Threshold + blur compute passes. ───────────────────────────────
+    std::array<UICustomBinding, 2> bloomBindings {{
+        { .group = 1, .binding = 0, .kind = UICustomBindingKind::SampledTexture, .pushOffset = 0 },
+        { .group = 1, .binding = 1, .kind = UICustomBindingKind::StorageTexture,
+          .format = static_cast<std::uint8_t>(WebGPUTexelFormat::RGBA16Float), .pushOffset = 0 },
+    }};
+    PlainComputeShader threshold;
+    threshold.Load(fs::path("threshold.comp.wgsl"), sizeof(Dim), bloomBindings);
+    PlainComputeShader blur;
+    blur.Load(fs::path("blur.comp.wgsl"), sizeof(Dim), bloomBindings);
+
+    std::array<std::uint32_t, 2> thresholdHandles { hdrScene.handle, bloomA.handle };
+    std::array<std::uint32_t, 2> blurHandles      { bloomA.handle,   bloomB.handle };
+
+    // ── Composite UI custom shader. ────────────────────────────────────
+    UIRenderer ui;
+    ui.Initialize(window, heap, cmd);
+
+    UICustomBinding compBindings[] = {
+        { .group = 2, .binding = 0, .kind = UICustomBindingKind::SampledTexture,
+          .pushOffset = static_cast<std::uint32_t>(offsetof(CompositePush, hdrSlot)) },
+        { .group = 2, .binding = 1, .kind = UICustomBindingKind::SampledTexture,
+          .pushOffset = static_cast<std::uint32_t>(offsetof(CompositePush, bloomSlot)) },
+        { .group = 2, .binding = 2, .kind = UICustomBindingKind::Sampler,
+          .pushOffset = static_cast<std::uint32_t>(offsetof(CompositePush, sampSlot)) },
+    };
+    WebGPUComputeShader composite;
+    composite.Load(fs::path("composite.comp.wgsl"), compBindings);
+
+    window.FinishInit();
+
+    // ── Passes: RT first (writes HDR scene), then UI (composite→canvas).
+    RTPass rtPass(&pipeline);
+    rtPass.handlesPtr   = rtHandles.data();
+    rtPass.handlesCount = static_cast<std::uint32_t>(rtHandles.size());
+    rtPass.maxDepth     = 1;                 // primary rays only
+    rtPass.outTexHandle = hdrScene.handle;   // ← RESOLVE target
+    window.passes.push_back(&rtPass);
+    window.passes.push_back(&ui);
+
+    // ── Static camera framing the grid from a front corner. ────────────
+    const float ext = float(kGrid - 1) * kSpacing;
+    Vector<float, 3, 4> camPos { ext * 1.1f, ext * 0.8f, ext * 1.6f + 3.0f };
+    Vector<float, 3, 4> d { -camPos.x, -camPos.y, -camPos.z };
+    const float dl = std::sqrt(d.x*d.x + d.y*d.y + d.z*d.z);
+    Vector<float, 3, 4> forward { d.x/dl, d.y/dl, d.z/dl };
+    Vector<float, 3, 4> worldUp { 0.0f, 1.0f, 0.0f };
+    Vector<float, 3, 4> right { forward.y*worldUp.z - forward.z*worldUp.y,
+                                forward.z*worldUp.x - forward.x*worldUp.z,
+                                forward.x*worldUp.y - forward.y*worldUp.x };
+    const float rl = std::sqrt(right.x*right.x + right.y*right.y + right.z*right.z);
+    right.x /= rl; right.y /= rl; right.z /= rl;
+    Vector<float, 3, 4> up { right.y*forward.z - right.z*forward.y,
+                             right.z*forward.x - right.x*forward.z,
+                             right.x*forward.y - right.y*forward.x };
+    {
+        CameraGPU& g = cameraBuf.value[0];
+        g.origin[0]=camPos.x; g.origin[1]=camPos.y; g.origin[2]=camPos.z; g.pad0=0;
+        g.right[0]=right.x;   g.right[1]=right.y;   g.right[2]=right.z;
+        g.up[0]=up.x;         g.up[1]=up.y;         g.up[2]=up.z;
+        g.forward[0]=forward.x; g.forward[1]=forward.y; g.forward[2]=forward.z;
+        g.aspect  = float(window.width) / float(window.height);
+        g.tanHalf = std::tan(60.0f * 3.14159265f / 360.0f);
+        g.pad1    = 0;
+        cameraBuf.FlushDevice();
+    }
+
+    // ── Bloom prepass: threshold + blur, each its own submit. ──────────
+    EventListener<void> bloomTick(&window.onBeforeUpdate, [&]() {
+        Dim dim { static_cast<std::uint32_t>(window.width),
+                  static_cast<std::uint32_t>(window.height), 0, 0 };
+        const std::uint32_t gx = (window.width  + 7u) / 8u;
+        const std::uint32_t gy = (window.height + 7u) / 8u;
+        threshold.Dispatch(&dim, sizeof(dim), thresholdHandles, gx, gy, 1);
+        blur.Dispatch(&dim, sizeof(dim), blurHandles, gx, gy, 1);
+    });
+
+    // ── Composite: scene + bloom → tonemap+gamma → canvas. ─────────────
+    EventListener<UIBuildArgs> composeSub(&ui.onBuild, [&](UIBuildArgs a) {
+        CompositePush pc { ui.FillHeader(0, 0), 0, 0, 0, 0 };
+        pc.hdrSlot   = static_cast<std::uint32_t>(static_cast<std::uint16_t>(hdrSlot));
+        pc.bloomSlot = static_cast<std::uint32_t>(static_cast<std::uint16_t>(bloomSlot));
+        pc.sampSlot  = static_cast<std::uint32_t>(static_cast<std::uint16_t>(sampSlot));
+        const std::uint32_t gx = (window.width  + 7u) / 8u;
+        const std::uint32_t gy = (window.height + 7u) / 8u;
+        ui.Dispatch(a.cmd, composite, &pc, sizeof(pc), gx, gy, 1);
+    });
+
+    std::println("[HDRBloom] RT→rgba16float→threshold→blur→composite running");
+
+    window.Render();
+    window.StartUpdate();
+    window.StartSync();
+    return 0;
+}
+#endif

examples/HDRBloom/miss.wgsl

@@ -0,0 +1,7 @@
+// HDRBloom miss (runs in SHADE). Dark, sub-threshold background so the
+// bloom pass only picks up the bright cubes, not the sky.
+fn miss_main(ray: RayDesc, payload: ptr<function, Payload>) {
+    let t = clamp(ray.direction.y * 0.5 + 0.5, 0.0, 1.0);
+    rtAccumulate(mix(vec3<f32>(0.015, 0.018, 0.030),
+                     vec3<f32>(0.030, 0.040, 0.070), t));
+}

examples/HDRBloom/project.cpp

@@ -0,0 +1,48 @@
+import std;
+import Crafter.Build;
+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;
+        }
+    }
+
+    std::vector<std::string> graphicsArgs(args.begin(), args.end());
+    Configuration* graphics = LocalProject({
+        .projectFile = "../../project.cpp",
+        .args = graphicsArgs,
+    });
+
+    Configuration cfg;
+    cfg.path = "./";
+    cfg.name = "HDRBloom";
+    cfg.outputName = "HDRBloom";
+    cfg.type = ConfigurationType::Executable;
+    if (isWasm) {
+        cfg.target = "wasm32-wasip1";
+        cfg.defines.push_back({"CRAFTER_GRAPHICS_WINDOW_DOM", ""});
+        cfg.compileFlags.push_back("-msimd128");
+    }
+    ApplyStandardArgs(cfg, args);
+    cfg.dependencies = { graphics };
+
+    std::array<fs::path, 0> ifaces = {};
+    std::array<fs::path, 1> impls = { "main" };
+    cfg.GetInterfacesAndImplementations(ifaces, impls);
+
+    if (isWasm) {
+        cfg.files.emplace_back(fs::path("raygen.wgsl"));
+        cfg.files.emplace_back(fs::path("closesthit.wgsl"));
+        cfg.files.emplace_back(fs::path("miss.wgsl"));
+        cfg.files.emplace_back(fs::path("threshold.comp.wgsl"));
+        cfg.files.emplace_back(fs::path("blur.comp.wgsl"));
+        cfg.files.emplace_back(fs::path("composite.comp.wgsl"));
+        EnableWasiBrowserRuntime(cfg);
+    }
+    return cfg;
+}

examples/HDRBloom/raygen.wgsl

@@ -0,0 +1,35 @@
+// HDRBloom raygen (runs in GENERATE). Host-driven pinhole camera at
+// @group(3) (groups 0..2 are reserved by the wavefront pipeline:
+// 0 = WfParams, 1 = data heaps, 2 = indirect args). Primary rays only —
+// maxDepth = 1.
+struct Camera {
+    origin:  vec3<f32>,
+    pad0:    f32,
+    right:   vec3<f32>,
+    tanHalf: f32,
+    up:      vec3<f32>,
+    aspect:  f32,
+    forward: vec3<f32>,
+    pad1:    f32,
+};
+@group(3) @binding(0) var<storage, read> camera : Camera;
+
+fn raygen_main(gid: vec3<u32>) {
+    if (gid.x >= wfParams.surfaceW || gid.y >= wfParams.surfaceH) { return; }
+
+    let pixelf = vec2<f32>(f32(gid.x), f32(gid.y));
+    let res    = vec2<f32>(f32(wfParams.surfaceW), f32(wfParams.surfaceH));
+    let uv     = (pixelf + vec2<f32>(0.5)) / res;
+    let ndc    = uv * 2.0 - vec2<f32>(1.0);
+
+    let direction = normalize(
+        camera.right   * (ndc.x  * camera.aspect * camera.tanHalf) +
+        camera.up      * (-ndc.y * camera.tanHalf) +
+        camera.forward);
+
+    var p: Payload;
+    p.color = vec3<f32>(0.0);
+
+    rtEmitPrimaryRay(camera.origin, 0.01, direction, 100000.0,
+                     0u, 0xFFu, 0u, 0u, p);
+}

examples/HDRBloom/threshold.comp.wgsl

@@ -0,0 +1,27 @@
+// HDRBloom threshold pass (PlainComputeShader). Reads the linear HDR scene
+// (rgba16float, written by the RT RESOLVE stage), keeps only the radiance
+// above 1.0, and writes it into the bloom mip (also rgba16float). This is
+// the primitive the issue asks for: sample a float texture (group 1
+// binding 0) and WRITE a user-owned float storage texture (group 1
+// binding 1) at an app-chosen binding.
+//
+// Layout (PlainComputeShader, no rayQuery → user groups start at 1):
+//   @group(0) @binding(0) uniform Dim         (surface size)
+//   @group(1) @binding(0) texture_2d<f32>     src   (sampled, float)
+//   @group(1) @binding(1) storage rgba16float dst   (write)
+
+struct Dim { w: u32, h: u32, _0: u32, _1: u32 };
+@group(0) @binding(0) var<uniform> dim : Dim;
+@group(1) @binding(0) var src : texture_2d<f32>;
+@group(1) @binding(1) var dst : texture_storage_2d<rgba16float, write>;
+
+@compute @workgroup_size(8, 8, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    if (gid.x >= dim.w || gid.y >= dim.h) { return; }
+    let coord = vec2<i32>(i32(gid.x), i32(gid.y));
+    let c = textureLoad(src, coord, 0).rgb;
+    // Soft knee around 1.0 so the extracted highlights ramp in smoothly.
+    let lum = dot(c, vec3<f32>(0.2126, 0.7152, 0.0722));
+    let keep = max(0.0, lum - 1.0) / max(lum, 1e-4);
+    textureStore(dst, coord, vec4<f32>(c * keep, 1.0));
+}

examples/README.md

@@ -90,3 +90,19 @@ Regression test for the WebGPU software ray-query shim. Builds a
 checking it against the analytically-known answer. Guards against the
 hardcoded-leaf-start TLAS-traversal bug (issue #25) that made every
 rayQuery pick miss for realistic instance counts. WebGPU/DOM only.
+
+### [HDRBloom](HDRBloom/)
+Cross-backend-shaped HDR bloom on the WebGPU backend (issue #27). The RT
+pipeline's RESOLVE writes linear radiance into a user `rgba16float`
+texture (`StorageImage2D`, `PipelineRTWebGPU::Init(..., RGBA16Float)`),
+two `PlainComputeShader` passes threshold + blur it through float storage
+targets (`UICustomBindingKind::StorageTexture`), and a UI custom shader
+composites scene + bloom with a Reinhard tonemap onto the canvas.
+Exercises the three primitives an HDR post-process chain needs: float
+textures, write-only storage-texture bindings, and pre-tonemap radiance
+out of the wavefront. The threshold/blur passes run from `onBeforeUpdate`
+so each gets its own queue submit — the storage-write → sampled-read
+barrier WebGPU only provides between submits (or between passes), never
+within a single compute pass. WebGPU/DOM only; the same chain is wireable
+on Vulkan today via an offscreen HDR heap image + a composite `RenderPass`
+(the present path records passes generically and barriers between them).