feat(webgpu): HDR post-process primitives — float textures, storage-texture bindings, RESOLVE→HDR target (#27)

Adds the minimal WebGPU-backend surface an HDR bloom/post chain needs:

1. wgpuCreateStorageImage2D + StorageImage2D<>: runtime rgba16float (or
   other format) textures with STORAGE_BINDING|TEXTURE_BINDING usage, no
   CompressedTextureAsset upload — GPU-produced HDR targets.
2. UICustomBindingKind::StorageTexture (kind 5): write-only storage-texture
   binding carrying a texel format (the freed UICustomBinding::format byte),
   wired through all three binding paths (UI custom / RT / plain compute).
   Float textures already sample via SampledTexture (sampleType float).
3. PipelineRTWebGPU::Init(..., hdrOutputFormat): RESOLVE writes the linear
   accumulator into a user rgba16float texture (RTPass::outTexHandle)
   instead of the rgba8unorm canvas, leaving the composite→swapchain pass to
   the app. Default RGBA8Unorm keeps the canvas path byte-identical.

Existing RT examples updated for the _pad→format field rename; the default
paths are unchanged (verified RTStress builds + renders).

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

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 }));

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;

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;

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)),

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 = 0, .kind = UICustomBindingKind::SampledTextureArray, .pushOffset = 0 },
-        { .group = 3, .binding = 1, .kind = UICustomBindingKind::Sampler,             ._pad = 0, .pushOffset = 0 },
+        { .group = 3, .binding = 1, .kind = UICustomBindingKind::Sampler,             .pushOffset = 0 },
-        { .group = 3, .binding = 2, .kind = UICustomBindingKind::Buffer,              ._pad = 0, .pushOffset = 0 },
+        { .group = 3, .binding = 2, .kind = UICustomBindingKind::Buffer,              .pushOffset = 0 },
     }};
 
     PipelineRTWebGPU pipeline;

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));
 }

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:

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:

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;

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);
         }
     };
 }

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

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);