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) and sample +// (texture_2d, 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;", + "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/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 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 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 hitGroups {{ { .type = RTShaderGroupType::TrianglesHitGroup, .closestHitShader = 2 } }}; std::array 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 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(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 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 missGroups, std::span hitGroups, const ShaderBindingTableWebGPU& sbt, - std::span bindings) { + std::span 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(wgsl.size()), userBindings.empty() ? nullptr : userBindings.data(), - static_cast(userBindings.size())); + static_cast(userBindings.size()), + static_cast(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 (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(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 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 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 raygenGroups, std::span missGroups, std::span hitGroups, const ShaderBindingTableWebGPU& sbt, - std::span bindings = {}); + std::span 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(gy), handlesPtr, static_cast(handlesCount), - static_cast(maxDepth)); + static_cast(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 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, 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 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);