diff --git a/WAVEFRONT-DESIGN.md b/WAVEFRONT-DESIGN.md
index 3b3742a..ff3ac7b 100644
--- a/WAVEFRONT-DESIGN.md
+++ b/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
diff --git a/additional/dom-webgpu.js b/additional/dom-webgpu.js
index facf92c..43971c8 100644
--- a/additional/dom-webgpu.js
+++ b/additional/dom-webgpu.js
@@ -46,6 +46,7 @@ function stub(name) {
         "wgpuCreateAtlasTexture", "wgpuWriteAtlasRegion", "wgpuDestroyTexture",
         "wgpuCreateImage2D", "wgpuWriteImage2D",
         "wgpuCreateImage2DArray", "wgpuWriteImage2DLayer",
+        "wgpuCreateStorageImage2D",
         "wgpuCreateLinearClampSampler", "wgpuCreateLinearRepeatSampler",
         "wgpuFrameBegin", "wgpuFrameEnd",
         "wgpuDispatchQuads", "wgpuDispatchCircles", "wgpuDispatchImages", "wgpuDispatchText",
@@ -835,6 +836,42 @@ env.wgpuWriteImage2DLayer = (handle, layer, level, srcPtr, byteSize, w, h) => {
     }
 };
 
+// Map the C++ WebGPUTexelFormat enum (u8) to a WGSL/GPU format string.
+// Kept deliberately small — the HDR post-process path needs rgba16float;
+// rgba8unorm (0) is the default so an unspecified `format` byte stays the
+// pre-existing behaviour. Storage-texture bindings (UICustomBindingKind 5)
+// and wgpuCreateStorageImage2D both run through here.
+function texelFormatStr(f) {
+    switch (f | 0) {
+        case 1:  return "rgba16float";
+        case 2:  return "r32float";
+        default: return "rgba8unorm";
+    }
+}
+
+// Runtime-allocated 2D texture with STORAGE_BINDING + TEXTURE_BINDING usage
+// and a caller-chosen format. Unlike wgpuCreateImage2D (rgba8unorm, sampled,
+// CPU-upload only) this is the allocation path an HDR post-process chain
+// needs: an rgba16float target a compute pass can BOTH write
+// (texture_storage_2d<rgba16float, write>) and sample
+// (texture_2d<f32>, filterable). No CompressedTextureAsset upload — the
+// content is produced on the GPU. COPY_DST so it can be cleared / seeded;
+// COPY_SRC so it can be read back or blitted.
+env.wgpuCreateStorageImage2D = (w, h, format) => {
+    const handle = newHandle();
+    const tex = device.createTexture({
+        size: [w, h],
+        format: texelFormatStr(format),
+        usage: GPUTextureUsage.STORAGE_BINDING
+             | GPUTextureUsage.TEXTURE_BINDING
+             | GPUTextureUsage.COPY_SRC
+             | GPUTextureUsage.COPY_DST,
+    });
+    textures.set(handle, tex);
+    textureViews.set(handle, tex.createView());
+    return handle;
+};
+
 env.wgpuWriteImage2D = (handle, srcPtr, byteSize, w, h) => {
     const tex = textures.get(handle);
     if (!tex) return;
@@ -1065,6 +1102,7 @@ env.wgpuLoadCustomShader = (wgslPtr, wgslLen, bindingsPtr, bindingsCount, rayQue
             group: dv.getUint8(i*8 + 0),
             binding: dv.getUint8(i*8 + 1),
             kind: dv.getUint8(i*8 + 2),
+            format: dv.getUint8(i*8 + 3),   // texel format for storage-texture kind
             pushOffset: dv.getUint32(i*8 + 4, true),
         });
     }
@@ -1130,6 +1168,7 @@ env.wgpuLoadCustomShader = (wgslPtr, wgslLen, bindingsPtr, bindingsCount, rayQue
                 else if (b.kind === 1) e.texture = { sampleType: "float", viewDimension: "2d" };
                 else if (b.kind === 2) e.sampler = { type: "filtering" };
                 else if (b.kind === 3) e.texture = { sampleType: "float", viewDimension: "2d-array" };
+                else if (b.kind === 5) e.storageTexture = { format: texelFormatStr(b.format), access: "write-only", viewDimension: "2d" };
                 return e;
             });
             bgls.push(device.createBindGroupLayout({ entries }));
@@ -1230,6 +1269,7 @@ env.wgpuDispatchCustom = (pipelineHandle, pushPtr, pushBytes, handlesPtr, handle
             else if (b.kind === 1) resource = textureViews.get(h);
             else if (b.kind === 2) resource = samplers.get(h);
             else if (b.kind === 3) resource = textureViews.get(h);
+            else if (b.kind === 5) resource = textureViews.get(h);
             return { binding: b.binding, resource };
         });
         const bg = device.createBindGroup({ layout: pipe.bgls[g], entries });
@@ -3060,9 +3100,17 @@ function wfLogTimestamps(ts, data) {
     console.log(`[crafter-wgpu] RT passes: ${parts.join(" | ")} | total ${(totalNs/1000).toFixed(1)}us`);
 }
 
-env.wgpuLoadRTPipeline = (wgslPtr, wgslLen, bindingsPtr, bindingsCount) => {
+env.wgpuLoadRTPipeline = (wgslPtr, wgslLen, bindingsPtr, bindingsCount, hdrOutputFormat) => {
     if (!rtState.vertHeap) rtInit();
     const userPart = new TextDecoder().decode(memU8().subarray(wgslPtr, wgslPtr + wgslLen));
+    // hdrOutputFormat (0 = none): when set, RESOLVE writes the linear
+    // accumulator into a user-provided storage texture of this format
+    // (rgba16float for HDR post-process) instead of the rgba8unorm canvas
+    // ping-pong. The app owns the downstream composite→swapchain pass. The
+    // binding-6 format must match in BOTH the WGSL `outImage` declaration
+    // and the data bind-group layout, so we swap both below.
+    const hdrOut = (hdrOutputFormat | 0) !== 0;
+    const outFmtStr = hdrOut ? texelFormatStr(hdrOutputFormat) : "rgba8unorm";
 
     // Insert helpers at the marker; prepend prelude.
     const marker = "// @CRAFTER_RT_LIBRARY_HELPERS_HERE";
@@ -3076,7 +3124,12 @@ env.wgpuLoadRTPipeline = (wgslPtr, wgslLen, bindingsPtr, bindingsCount) => {
     // Wavefront assembly: types + bindings | user CH/miss/resolve + wfPayload
     // + switches (beforeHelpers) | pure helpers | wavefront helpers | user
     // raygen + the five @compute entry points (afterHelpers).
-    const fullWgsl = rtWgslTypes + rtWgslWavefrontBindings + "\n"
+    const wavefrontBindings = hdrOut
+        ? rtWgslWavefrontBindings.replace(
+              "var outImage : texture_storage_2d<rgba8unorm, write>;",
+              "var outImage : texture_storage_2d<" + outFmtStr + ", write>;")
+        : rtWgslWavefrontBindings;
+    const fullWgsl = rtWgslTypes + wavefrontBindings + "\n"
         + beforeHelpers + "\n" + rtWgslPureHelpers + "\n"
         + rtWgslWavefrontHelpers + "\n" + afterHelpers;
 
@@ -3102,6 +3155,7 @@ env.wgpuLoadRTPipeline = (wgslPtr, wgslLen, bindingsPtr, bindingsCount) => {
                 group:      g,
                 binding:    dv.getUint8(i*8 + 1),
                 kind:       dv.getUint8(i*8 + 2),
+                format:     dv.getUint8(i*8 + 3),   // texel format for storage-texture kind
                 pushOffset: dv.getUint32(i*8 + 4, true),
             });
         }
@@ -3124,7 +3178,7 @@ env.wgpuLoadRTPipeline = (wgslPtr, wgslLen, bindingsPtr, bindingsCount) => {
         const dataBgl = device.createBindGroupLayout({ entries: [
             sb(0), sb(1), sb(2), sb(3), sb(4), sb(5),
             { binding: 6, visibility: GPUShaderStage.COMPUTE,
-              storageTexture: { format: "rgba8unorm", access: "write-only", viewDimension: "2d" } },
+              storageTexture: { format: outFmtStr, access: "write-only", viewDimension: "2d" } },
             sb(7), sb(8), sb(9),
             rw(10), rw(11), rw(12), rw(13), rw(14), rw(15),
         ]});
@@ -3143,6 +3197,7 @@ env.wgpuLoadRTPipeline = (wgslPtr, wgslLen, bindingsPtr, bindingsCount) => {
                     else if (b.kind === 2) e.sampler = { type: "filtering" };
                     else if (b.kind === 3) e.texture = { sampleType: "float", viewDimension: "2d-array" };
                     else if (b.kind === 4) e.buffer  = { type: "storage" };
+                    else if (b.kind === 5) e.storageTexture = { format: texelFormatStr(b.format), access: "write-only", viewDimension: "2d" };
                     return e;
                 });
                 userBgls.push(device.createBindGroupLayout({ entries }));
@@ -3172,7 +3227,7 @@ env.wgpuLoadRTPipeline = (wgslPtr, wgslLen, bindingsPtr, bindingsCount) => {
             shadePipe:   mk(userLayout,  "wfShade"),
             resolvePipe: mk(userLayout,  "wfResolve"),
             paramsBgl, dataBgl, indirectBgl, emptyBgl, userBgls,
-            byGroup, sortedGroups, traceHasUser,
+            byGroup, sortedGroups, traceHasUser, hdrOutput: hdrOut,
         };
         const handle = newHandle();
         rtPipelines.set(handle, entry);
@@ -3203,6 +3258,7 @@ function wfUserBindGroups(pipe, handlesPtr, handlesCount) {
                 else if (b.kind === 2) resource = samplers.get(h);
                 else if (b.kind === 3) resource = textureViews.get(h);
                 else if (b.kind === 4) resource = { buffer: buffers.get(h) };
+                else if (b.kind === 5) resource = textureViews.get(h);
                 return { binding: b.binding, resource };
             });
             out.push({ group: g, bindGroup: device.createBindGroup({ layout: pipe.userBgls[bglIdx], entries }) });
@@ -3214,7 +3270,7 @@ function wfUserBindGroups(pipe, handlesPtr, handlesCount) {
 
 env.wgpuDispatchRT = (pipelineHandle, pushPtr, pushBytes,
                       tlasBufHandle, instanceCount, gx, gy,
-                      handlesPtr, handlesCount, maxDepth) => {
+                      handlesPtr, handlesCount, maxDepth, outTexHandle) => {
     if (!state.encoder) return;
     const pipe = rtPipelines.get(pipelineHandle);
     const tlas = buffers.get(tlasBufHandle);
@@ -3222,6 +3278,17 @@ env.wgpuDispatchRT = (pipelineHandle, pushPtr, pushBytes,
         console.error("[crafter-wgpu] wgpuDispatchRT: unknown pipeline or tlas");
         return;
     }
+    // HDR-output pipelines write RESOLVE into a user-provided storage
+    // texture (binding 6) rather than the canvas ping-pong. The handle is
+    // required in that case and its format must match the one the pipeline
+    // was loaded with.
+    if (pipe.hdrOutput) {
+        const userOut = textureViews.get(outTexHandle >>> 0);
+        if (!userOut) {
+            console.error("[crafter-wgpu] wgpuDispatchRT: HDR pipeline needs a valid outTexHandle");
+            return;
+        }
+    }
     const entryOrderBuf = buffers.get(rtState.currentEntryOrder);
     const bvhBuf        = buffers.get(rtState.currentBvh);
     if (!entryOrderBuf || !bvhBuf) {
@@ -3259,7 +3326,9 @@ env.wgpuDispatchRT = (pipelineHandle, pushPtr, pushBytes,
     writeSlot(1 + 3 * depth, 1, depth);              // RESOLVE
     queue.writeBuffer(wf.paramsRing, 0, ring, 0, passCount * 64);
 
-    const outView = state.outIsPing ? state.pingView : state.pongView;
+    const outView = pipe.hdrOutput
+        ? textureViews.get(outTexHandle >>> 0)
+        : (state.outIsPing ? state.pingView : state.pongView);
     const paramsBg = device.createBindGroup({
         layout: pipe.paramsBgl,
         entries: [{ binding: 0, resource: { buffer: wf.paramsRing, offset: 0, size: 256 } }],
@@ -3389,9 +3458,12 @@ env.wgpuDispatchRT = (pipelineHandle, pushPtr, pushBytes,
     }
 
     // Reopen the frame's shared pass so wgpuFrameEnd / later UI work as
-    // before, and flip ping-pong so the blit picks the texture RESOLVE wrote.
+    // before. For canvas output, flip ping-pong so the blit picks the
+    // texture RESOLVE wrote. For HDR output RESOLVE wrote a user texture and
+    // left the ping-pong untouched, so the flip is skipped — the app's own
+    // composite pass is responsible for writing the canvas.
     state.pass = enc.beginComputePass();
-    state.outIsPing = !state.outIsPing;
+    if (!pipe.hdrOutput) state.outIsPing = !state.outIsPing;
 };
 
 // ── Standalone compute pipelines ────────────────────────────────────────
@@ -3434,6 +3506,7 @@ env.wgpuLoadComputePipeline = (wgslPtr, wgslLen, pushUniformSize,
                 group:   dv.getUint8(i*8 + 0),
                 binding: dv.getUint8(i*8 + 1),
                 kind:    dv.getUint8(i*8 + 2),
+                format:  dv.getUint8(i*8 + 3),   // texel format for storage-texture kind
                 pushOffset: dv.getUint32(i*8 + 4, true),
             });
         }
@@ -3492,6 +3565,7 @@ env.wgpuLoadComputePipeline = (wgslPtr, wgslLen, pushUniformSize,
                 else if (b.kind === 2) e.sampler = { type: "filtering" };
                 else if (b.kind === 3) e.texture = { sampleType: "float", viewDimension: "2d-array" };
                 else if (b.kind === 4) e.buffer  = { type: "storage" };  // read-write storage
+                else if (b.kind === 5) e.storageTexture = { format: texelFormatStr(b.format), access: "write-only", viewDimension: "2d" };  // write-only storage texture
                 return e;
             });
             bgls.push(device.createBindGroupLayout({ entries }));
@@ -3620,6 +3694,7 @@ env.wgpuDispatchCompute = (pipelineHandle, pushPtr, pushBytes,
                     else if (b.kind === 1)                 resource = textureViews.get(h);
                     else if (b.kind === 2)                 resource = samplers.get(h);
                     else if (b.kind === 3)                 resource = textureViews.get(h);
+                    else if (b.kind === 5)                 resource = textureViews.get(h);
                     return { binding: b.binding, resource };
                 });
                 userBGs.push(device.createBindGroup({ layout: pipe.bgls[bglIdx++], entries }));
diff --git a/examples/HDRBloom/blur.comp.wgsl b/examples/HDRBloom/blur.comp.wgsl
new file mode 100644
index 0000000..b4a570a
--- /dev/null
+++ b/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));
+}
diff --git a/examples/HDRBloom/closesthit.wgsl b/examples/HDRBloom/closesthit.wgsl
new file mode 100644
index 0000000..ef83cf8
--- /dev/null
+++ b/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)));
+    }
+}
diff --git a/examples/HDRBloom/composite.comp.wgsl b/examples/HDRBloom/composite.comp.wgsl
new file mode 100644
index 0000000..f1c893a
--- /dev/null
+++ b/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));
+}
diff --git a/examples/HDRBloom/main.cpp b/examples/HDRBloom/main.cpp
new file mode 100644
index 0000000..6495c19
--- /dev/null
+++ b/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
diff --git a/examples/HDRBloom/miss.wgsl b/examples/HDRBloom/miss.wgsl
new file mode 100644
index 0000000..fcaae4f
--- /dev/null
+++ b/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));
+}
diff --git a/examples/HDRBloom/project.cpp b/examples/HDRBloom/project.cpp
new file mode 100644
index 0000000..071adba
--- /dev/null
+++ b/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;
+}
diff --git a/examples/HDRBloom/raygen.wgsl b/examples/HDRBloom/raygen.wgsl
new file mode 100644
index 0000000..ad45486
--- /dev/null
+++ b/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);
+}
diff --git a/examples/HDRBloom/threshold.comp.wgsl b/examples/HDRBloom/threshold.comp.wgsl
new file mode 100644
index 0000000..8003f86
--- /dev/null
+++ b/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));
+}
diff --git a/examples/README.md b/examples/README.md
index 940d54d..ed4cb2a 100644
--- a/examples/README.md
+++ b/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).
diff --git a/examples/RTStress/main.cpp b/examples/RTStress/main.cpp
index d9581a2..fb64d4e 100644
--- a/examples/RTStress/main.cpp
+++ b/examples/RTStress/main.cpp
@@ -62,7 +62,7 @@ int main() {
 
     // One user binding: the camera storage buffer at @group(3).
     std::array<UICustomBinding, 1> bindings {{
-        { .group = 3, .binding = 0, .kind = UICustomBindingKind::Buffer, ._pad = 0, .pushOffset = 0 },
+        { .group = 3, .binding = 0, .kind = UICustomBindingKind::Buffer, .pushOffset = 0 },
     }};
 
     PipelineRTWebGPU pipeline;
diff --git a/examples/RTVolume/main.cpp b/examples/RTVolume/main.cpp
index 48cbed2..fc3ff69 100644
--- a/examples/RTVolume/main.cpp
+++ b/examples/RTVolume/main.cpp
@@ -74,7 +74,7 @@ int main() {
     } }};
 
     std::array<UICustomBinding, 1> bindings {{
-        { .group = 3, .binding = 0, .kind = UICustomBindingKind::Buffer, ._pad = 0, .pushOffset = 0 },
+        { .group = 3, .binding = 0, .kind = UICustomBindingKind::Buffer, .pushOffset = 0 },
     }};
 
     PipelineRTWebGPU pipeline;
diff --git a/examples/RayQueryPick/main.cpp b/examples/RayQueryPick/main.cpp
index 157282b..4cb0606 100644
--- a/examples/RayQueryPick/main.cpp
+++ b/examples/RayQueryPick/main.cpp
@@ -87,7 +87,7 @@ int main() {
     std::array<RTShaderGroup, 1> hitGroups    {{ { .type = RTShaderGroupType::TrianglesHitGroup, .closestHitShader = 2 } }};
 
     std::array<UICustomBinding, 1> bindings {{
-        { .group = 3, .binding = 0, .kind = UICustomBindingKind::Buffer, ._pad = 0, .pushOffset = 0 },
+        { .group = 3, .binding = 0, .kind = UICustomBindingKind::Buffer, .pushOffset = 0 },
     }};
 
     PipelineRTWebGPU pipeline;
@@ -172,7 +172,7 @@ int main() {
     // ── rayQuery pick shader + output buffer. ──────────────────────────
     static PlainComputeShader pickShader;
     std::array<UICustomBinding, 1> pickBindings {{
-        { .group = 2, .binding = 0, .kind = UICustomBindingKind::BufferReadWrite, ._pad = 0, .pushOffset = 0 },
+        { .group = 2, .binding = 0, .kind = UICustomBindingKind::BufferReadWrite, .pushOffset = 0 },
     }};
     pickShader.Load(fs::path("rayquery_pick.wgsl"),
                     static_cast<std::uint32_t>(sizeof(PickPush)),
diff --git a/examples/Sponza/main.cpp b/examples/Sponza/main.cpp
index 1df3a2c..0d4c265 100644
--- a/examples/Sponza/main.cpp
+++ b/examples/Sponza/main.cpp
@@ -278,9 +278,9 @@ int main() {
     //   binding 1 — sampler (linear clamp)
     //   binding 2 — Camera storage buffer (host-driven, updated per frame)
     std::array<UICustomBinding, 3> bindings {{
-        { .group = 3, .binding = 0, .kind = UICustomBindingKind::SampledTextureArray, ._pad = 0, .pushOffset = 0 },
-        { .group = 3, .binding = 1, .kind = UICustomBindingKind::Sampler,             ._pad = 0, .pushOffset = 0 },
-        { .group = 3, .binding = 2, .kind = UICustomBindingKind::Buffer,              ._pad = 0, .pushOffset = 0 },
+        { .group = 3, .binding = 0, .kind = UICustomBindingKind::SampledTextureArray, .pushOffset = 0 },
+        { .group = 3, .binding = 1, .kind = UICustomBindingKind::Sampler,             .pushOffset = 0 },
+        { .group = 3, .binding = 2, .kind = UICustomBindingKind::Buffer,              .pushOffset = 0 },
     }};
 
     PipelineRTWebGPU pipeline;
diff --git a/implementations/Crafter.Graphics-PipelineRTWebGPU.cpp b/implementations/Crafter.Graphics-PipelineRTWebGPU.cpp
index 27642a5..2de9cf8 100644
--- a/implementations/Crafter.Graphics-PipelineRTWebGPU.cpp
+++ b/implementations/Crafter.Graphics-PipelineRTWebGPU.cpp
@@ -80,8 +80,10 @@ void PipelineRTWebGPU::Init(WebGPUCommandEncoderRef                 /*cmd*/,
                             std::span<const RTShaderGroup>          missGroups,
                             std::span<const RTShaderGroup>          hitGroups,
                             const ShaderBindingTableWebGPU&         sbt,
-                            std::span<const UICustomBinding>        bindings) {
+                            std::span<const UICustomBinding>        bindings,
+                            WebGPUTexelFormat                       hdrOutputFormat) {
     userBindings.assign(bindings.begin(), bindings.end());
+    hdrOutput = (hdrOutputFormat != WebGPUTexelFormat::RGBA8Unorm);
     std::string wgsl;
     wgsl.reserve(8 * 1024);
 
@@ -306,5 +308,6 @@ void PipelineRTWebGPU::Init(WebGPUCommandEncoderRef                 /*cmd*/,
         wgsl.data(),
         static_cast<std::int32_t>(wgsl.size()),
         userBindings.empty() ? nullptr : userBindings.data(),
-        static_cast<std::int32_t>(userBindings.size()));
+        static_cast<std::int32_t>(userBindings.size()),
+        static_cast<std::int32_t>(hdrOutputFormat));
 }
diff --git a/implementations/Crafter.Graphics-UI-WebGPU.cpp b/implementations/Crafter.Graphics-UI-WebGPU.cpp
index 2c46c7c..ef3d90e 100644
--- a/implementations/Crafter.Graphics-UI-WebGPU.cpp
+++ b/implementations/Crafter.Graphics-UI-WebGPU.cpp
@@ -174,6 +174,8 @@ void UIRenderer::Dispatch(GraphicsCommandBuffer /*cmd*/, const GraphicsComputeSh
                 if (slot < heap_->bufferTable.size()) handle = heap_->bufferTable[slot];
                 break;
             case UICustomBindingKind::SampledTexture:
+            case UICustomBindingKind::SampledTextureArray:
+            case UICustomBindingKind::StorageTexture:
                 if (slot < heap_->imageTable.size()) handle = heap_->imageTable[slot];
                 break;
             case UICustomBindingKind::Sampler:
diff --git a/interfaces/Crafter.Graphics-Image2D.cppm b/interfaces/Crafter.Graphics-Image2D.cppm
index 369d539..07d09be 100644
--- a/interfaces/Crafter.Graphics-Image2D.cppm
+++ b/interfaces/Crafter.Graphics-Image2D.cppm
@@ -47,8 +47,48 @@ import std;
 import Crafter.Asset;
 import :DescriptorHeapWebGPU;
 import :WebGPU;
+import :WebGPUComputeShader;   // WebGPUTexelFormat
 
 export namespace Crafter {
+    // GPU-produced 2D texture with STORAGE_BINDING + TEXTURE_BINDING usage
+    // and a runtime-chosen format. The HDR post-process building block: an
+    // rgba16float target a compute pass can write through a StorageTexture
+    // binding and a later pass can sample through a SampledTexture binding.
+    // No CompressedTextureAsset / CPU upload — content comes from the GPU.
+    // Distinct from Image2D<T> (rgba8unorm, sampled-only, asset-upload).
+    class StorageImage2D {
+    public:
+        WebGPUTextureRef  handle = 0;
+        std::uint16_t     width  = 0;
+        std::uint16_t     height = 0;
+        WebGPUTexelFormat format = WebGPUTexelFormat::RGBA16Float;
+
+        void Create(std::uint16_t w, std::uint16_t h,
+                    WebGPUTexelFormat fmt = WebGPUTexelFormat::RGBA16Float) {
+            width  = w;
+            height = h;
+            format = fmt;
+            handle = WebGPU::wgpuCreateStorageImage2D(
+                w, h, static_cast<std::int32_t>(fmt));
+        }
+
+        // Register in a heap image slot so a UI custom shader can bind it
+        // via UICustomBinding (SampledTexture to read, StorageTexture to
+        // write). PlainComputeShader / RTPass take the raw `handle` directly.
+        ImageSlot AllocateSlot(DescriptorHeapWebGPU& heap) {
+            DescriptorRange r = heap.AllocateImageSlots(1);
+            heap.imageTable[r.firstElement] = handle;
+            return ImageSlot(&heap, r.firstElement);
+        }
+
+        void Destroy() {
+            if (handle != 0) {
+                WebGPU::wgpuDestroyTexture(handle);
+                handle = 0;
+            }
+        }
+    };
+
     template <typename PixelType>
     class Image2D {
     public:
diff --git a/interfaces/Crafter.Graphics-PipelineRTWebGPU.cppm b/interfaces/Crafter.Graphics-PipelineRTWebGPU.cppm
index 6825fbd..a286994 100644
--- a/interfaces/Crafter.Graphics-PipelineRTWebGPU.cppm
+++ b/interfaces/Crafter.Graphics-PipelineRTWebGPU.cppm
@@ -36,6 +36,10 @@ export namespace Crafter {
         // RTPass to consult when packing the handles[] array at dispatch
         // time (one resolved u32 handle per binding, in declaration order).
         std::vector<UICustomBinding> userBindings;
+        // True when Init was given a non-default hdrOutputFormat: RESOLVE
+        // writes a user storage texture instead of the canvas, and RTPass
+        // must supply its handle via RTPass::outTexHandle.
+        bool hdrOutput = false;
 
         // Build the megakernel pipeline. Groups carry indices into
         // `sbt.shaders`. The library generates one `case` per registered
@@ -45,12 +49,19 @@ export namespace Crafter {
         // `userBindings` declares extra @group(2)+ resources the user's
         // closest-hit / miss / raygen WGSL touches (material SSBOs,
         // albedo textures, samplers).
+        // `hdrOutputFormat` defaults to RGBA8Unorm — the canvas ping-pong
+        // path, unchanged. Pass RGBA16Float (or another float format) to
+        // have RESOLVE write the linear accumulator into a user storage
+        // texture instead, for an HDR post-process chain. The app then owns
+        // the composite→swapchain pass and must set RTPass::outTexHandle.
         void Init(WebGPUCommandEncoderRef                 cmd,
                   std::span<const RTShaderGroup>          raygenGroups,
                   std::span<const RTShaderGroup>          missGroups,
                   std::span<const RTShaderGroup>          hitGroups,
                   const ShaderBindingTableWebGPU&         sbt,
-                  std::span<const UICustomBinding>        bindings = {});
+                  std::span<const UICustomBinding>        bindings = {},
+                  WebGPUTexelFormat                       hdrOutputFormat
+                      = WebGPUTexelFormat::RGBA8Unorm);
 
         PipelineRTWebGPU() = default;
         PipelineRTWebGPU(const PipelineRTWebGPU&) = delete;
diff --git a/interfaces/Crafter.Graphics-RTPass.cppm b/interfaces/Crafter.Graphics-RTPass.cppm
index f03f2c9..75a7edd 100644
--- a/interfaces/Crafter.Graphics-RTPass.cppm
+++ b/interfaces/Crafter.Graphics-RTPass.cppm
@@ -97,6 +97,12 @@ export namespace Crafter {
         // bounce; etc. The library unrolls GENERATE; (PREP; TRACE; SHADE)
         // ×maxDepth; RESOLVE.
         std::uint32_t     maxDepth     = 1;
+        // Destination storage-texture handle for an HDR-output pipeline
+        // (one Init'd with a non-default hdrOutputFormat). RESOLVE writes
+        // the linear accumulator here instead of the canvas; the app runs
+        // its own composite→swapchain pass afterwards. Ignored (0) for the
+        // default canvas path.
+        std::uint32_t     outTexHandle = 0;
 
         RTPass(PipelineRTWebGPU* p) : pipeline(p) {}
 
@@ -114,7 +120,8 @@ export namespace Crafter {
                 static_cast<std::int32_t>(gy),
                 handlesPtr,
                 static_cast<std::int32_t>(handlesCount),
-                static_cast<std::int32_t>(maxDepth));
+                static_cast<std::int32_t>(maxDepth),
+                outTexHandle);
         }
     };
 }
diff --git a/interfaces/Crafter.Graphics-WebGPU.cppm b/interfaces/Crafter.Graphics-WebGPU.cppm
index 2122cba..cffc122 100644
--- a/interfaces/Crafter.Graphics-WebGPU.cppm
+++ b/interfaces/Crafter.Graphics-WebGPU.cppm
@@ -109,6 +109,17 @@ namespace Crafter::WebGPU {
                                           const void* srcPtr, std::int32_t byteSize,
                                           std::int32_t w, std::int32_t h);
 
+    // Runtime storage texture — STORAGE_BINDING | TEXTURE_BINDING usage,
+    // caller-chosen format (`format` is the WebGPUTexelFormat enum, e.g.
+    // 1 = rgba16float). Unlike wgpuCreateImage2D this needs no
+    // CompressedTextureAsset upload: the content is produced on the GPU
+    // (e.g. an HDR post-process target a compute pass writes via a
+    // StorageTexture binding and later samples via SampledTexture). Backs
+    // StorageImage2D<T> in :Image2D.
+    __attribute__((import_module("env"), import_name("wgpuCreateStorageImage2D")))
+    extern "C" std::uint32_t wgpuCreateStorageImage2D(std::int32_t w, std::int32_t h,
+                                                      std::int32_t format);
+
     __attribute__((import_module("env"), import_name("wgpuCreateLinearClampSampler")))
     extern "C" std::uint32_t wgpuCreateLinearClampSampler();
 
@@ -190,10 +201,16 @@ namespace Crafter::WebGPU {
     // UICustomBinding-shaped (8 bytes each) declaring extra @group(2)+
     // resources the user's closest-hit / miss / raygen WGSL references.
     // Pass (nullptr, 0) for a pipeline with no user-declared bindings.
+    // `hdrOutputFormat` (0 = none) makes the RESOLVE stage write the linear
+    // accumulator into a user-supplied storage texture of that
+    // WebGPUTexelFormat (e.g. 1 = rgba16float) instead of the rgba8unorm
+    // canvas ping-pong — the entry point for an HDR post-process chain. The
+    // matching output texture handle is passed to wgpuDispatchRT.
     // Returns an opaque pipeline handle.
     __attribute__((import_module("env"), import_name("wgpuLoadRTPipeline")))
     extern "C" std::uint32_t wgpuLoadRTPipeline(const void* wgslPtr, std::int32_t wgslLen,
-                                                const void* bindingsPtr, std::int32_t bindingsCount);
+                                                const void* bindingsPtr, std::int32_t bindingsCount,
+                                                std::int32_t hdrOutputFormat);
 
     // Dispatch a TraceRays-equivalent pass: the RT pipeline is dispatched
     // over a (gx, gy) tile grid; the library writes the push data (camera,
@@ -202,6 +219,9 @@ namespace Crafter::WebGPU {
     // `handles[]` carries resolved WebGPU resource handles for every user
     // binding declared at pipeline-load time, in the same order. Pass
     // (nullptr, 0) for a pipeline with no user bindings.
+    // `outTexHandle` is the destination texture for an HDR-output pipeline
+    // (one loaded with hdrOutputFormat != 0); ignored (pass 0) for the
+    // default canvas path.
     __attribute__((import_module("env"), import_name("wgpuDispatchRT")))
     extern "C" void wgpuDispatchRT(std::uint32_t pipelineHandle,
                                    const void* pushPtr, std::int32_t pushBytes,
@@ -209,7 +229,8 @@ namespace Crafter::WebGPU {
                                    std::int32_t  instanceCount,
                                    std::int32_t  gx, std::int32_t gy,
                                    const void* handlesPtr, std::int32_t handlesCount,
-                                   std::int32_t maxDepth);
+                                   std::int32_t maxDepth,
+                                   std::uint32_t outTexHandle);
 
     // GPU TLAS-build dispatch. Two sequential compute passes:
     //   1. tlasBuildMain — per-instance world AABB + identity permutation
diff --git a/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm b/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm
index 68a4983..58f0f77 100644
--- a/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm
+++ b/interfaces/Crafter.Graphics-WebGPUComputeShader.cppm
@@ -31,6 +31,16 @@ import std;
 import :WebGPU;
 
 export namespace Crafter {
+    // Texel format for runtime-allocated textures and write-only storage-
+    // texture bindings. Values are the wire contract with dom-webgpu.js's
+    // texelFormatStr(); keep them in sync. rgba8unorm (0) is the default so
+    // an unset `format` byte preserves prior behaviour.
+    enum class WebGPUTexelFormat : std::uint8_t {
+        RGBA8Unorm  = 0,
+        RGBA16Float = 1,   // HDR post-process targets — filterable per spec
+        R32Float    = 2,
+    };
+
     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
@@ -41,13 +51,20 @@ export namespace Crafter {
         // integrate node momentum, write brace stress, or output TLAS
         // instance transforms.
         BufferReadWrite     = 4,
+        // write-only storage texture (texture_storage_2d<FMT, write> in
+        // WGSL). The `format` field carries the WebGPUTexelFormat. Lets a
+        // compute pass write a user-owned texture (e.g. an rgba16float bloom
+        // mip) at an arbitrary group/binding — the missing primitive for an
+        // HDR mip pyramid. Bind the same texture as SampledTexture to read
+        // it back in a later pass.
+        StorageTexture      = 5,
     };
 
     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::uint8_t  format;       // WebGPUTexelFormat for StorageTexture; 0 otherwise
         std::uint32_t pushOffset;   // offset in push data where the slot uint32 lives
     };
     static_assert(sizeof(UICustomBinding) == 8);