From bbe1b21c22f9a4c8235377e6acd5c85b392d2ada Mon Sep 17 00:00:00 2001 From: catbot Date: Tue, 9 Jun 2026 12:55:14 +0000 Subject: [PATCH] test(webgpu): HDRBloom example exercising the HDR post-process primitives (#27) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit RT→rgba16float→threshold→blur→composite, end-to-end proof of the three primitives. threshold/blur run from onBeforeUpdate (one submit each) so the storage-write→sampled-read dependency gets WebGPU's per-submit barrier (there is none between dispatches within a compute pass); the composite is a UI custom shader so it owns the canvas ping-pong. Documents the Vulkan symmetry (gap 4): the native present path records passes generically and barriers between them, so the same chain is wireable today. Co-Authored-By: Claude Opus 4.8 --- WAVEFRONT-DESIGN.md | 10 ++ examples/HDRBloom/blur.comp.wgsl | 33 ++++ examples/HDRBloom/closesthit.wgsl | 45 +++++ examples/HDRBloom/composite.comp.wgsl | 53 ++++++ examples/HDRBloom/main.cpp | 244 ++++++++++++++++++++++++++ examples/HDRBloom/miss.wgsl | 7 + examples/HDRBloom/project.cpp | 48 +++++ examples/HDRBloom/raygen.wgsl | 35 ++++ examples/HDRBloom/threshold.comp.wgsl | 27 +++ examples/README.md | 16 ++ 10 files changed, 518 insertions(+) create mode 100644 examples/HDRBloom/blur.comp.wgsl create mode 100644 examples/HDRBloom/closesthit.wgsl create mode 100644 examples/HDRBloom/composite.comp.wgsl create mode 100644 examples/HDRBloom/main.cpp create mode 100644 examples/HDRBloom/miss.wgsl create mode 100644 examples/HDRBloom/project.cpp create mode 100644 examples/HDRBloom/raygen.wgsl create mode 100644 examples/HDRBloom/threshold.comp.wgsl 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, hdr:vec4)->vec4`. 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/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 dim : Dim; +@group(1) @binding(0) var src : texture_2d; +@group(1) @binding(1) var dst : texture_storage_2d; + +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) { + 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(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(x, y), 0).rgb; + count = count + 1.0; + } + } + textureStore(dst, vec2(i32(gid.x), i32(gid.y)), vec4(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, +}; + +const SUN_DIR_TO_LIGHT: vec3 = vec3(0.40, 0.85, 0.35); +const AMBIENT_COLOR: vec3 = vec3(0.10, 0.11, 0.16); + +// Distinct per-instance hue so emitters read as different coloured glows. +fn instanceAlbedo(i: u32) -> vec3 { + let h = i * 2654435761u; + return vec3( + 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) { + 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( + 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(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 out (auto) +// group 1 binding 1 — texture_2d prev (auto, unused) +// group 2 binding 0 — texture_2d hdr scene (heap slot in push) +// group 2 binding 1 — texture_2d 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 hdr : UIDispatchHeader; +@group(1) @binding(0) var outTex : texture_storage_2d; +@group(1) @binding(1) var prevTex : texture_2d; +@group(2) @binding(0) var hdrTex : texture_2d; +@group(2) @binding(1) var bloomTex : texture_2d; +@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) { + if (gid.x >= hdr.surfaceW || gid.y >= hdr.surfaceH) { return; } + let coord = vec2(i32(gid.x), i32(gid.y)); + + let scene = textureLoad(hdrTex, coord, 0).rgb; + let uv = (vec2(f32(gid.x), f32(gid.y)) + vec2(0.5)) + / vec2(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(1.0)); // Reinhard tonemap + let g = pow(mapped, vec3(1.0 / 2.2)); // gamma 2.2 + textureStore(outTex, coord, vec4(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 // 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(window.width); + const std::uint16_t H = static_cast(window.height); + + // ── RT pipeline: HDR output (RESOLVE → rgba16float) ──────────────── + std::array 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 raygenGroups {{ { .type = RTShaderGroupType::General, .generalShader = 0 } }}; + std::array missGroups {{ { .type = RTShaderGroupType::General, .generalShader = 1 } }}; + std::array hitGroups {{ { .type = RTShaderGroupType::TrianglesHitGroup, .closestHitShader = 2 } }}; + + std::array 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, 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 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 cameraBuf; + cameraBuf.Create(1); + static std::array rtHandles { cameraBuf.handle }; + + static std::vector renderers; + renderers.reserve(static_cast(kGrid * kGrid * kGrid)); + const float origin0 = -0.5f * static_cast(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(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 bloomBindings {{ + { .group = 1, .binding = 0, .kind = UICustomBindingKind::SampledTexture, .pushOffset = 0 }, + { .group = 1, .binding = 1, .kind = UICustomBindingKind::StorageTexture, + .format = static_cast(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 thresholdHandles { hdrScene.handle, bloomA.handle }; + std::array 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(offsetof(CompositePush, hdrSlot)) }, + { .group = 2, .binding = 1, .kind = UICustomBindingKind::SampledTexture, + .pushOffset = static_cast(offsetof(CompositePush, bloomSlot)) }, + { .group = 2, .binding = 2, .kind = UICustomBindingKind::Sampler, + .pushOffset = static_cast(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(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 camPos { ext * 1.1f, ext * 0.8f, ext * 1.6f + 3.0f }; + Vector 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 forward { d.x/dl, d.y/dl, d.z/dl }; + Vector worldUp { 0.0f, 1.0f, 0.0f }; + Vector 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 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 bloomTick(&window.onBeforeUpdate, [&]() { + Dim dim { static_cast(window.width), + static_cast(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 composeSub(&ui.onBuild, [&](UIBuildArgs a) { + CompositePush pc { ui.FillHeader(0, 0), 0, 0, 0, 0 }; + pc.hdrSlot = static_cast(static_cast(hdrSlot)); + pc.bloomSlot = static_cast(static_cast(bloomSlot)); + pc.sampSlot = static_cast(static_cast(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) { + let t = clamp(ray.direction.y * 0.5 + 0.5, 0.0, 1.0); + rtAccumulate(mix(vec3(0.015, 0.018, 0.030), + vec3(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 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 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 ifaces = {}; + std::array 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, + pad0: f32, + right: vec3, + tanHalf: f32, + up: vec3, + aspect: f32, + forward: vec3, + pad1: f32, +}; +@group(3) @binding(0) var camera : Camera; + +fn raygen_main(gid: vec3) { + if (gid.x >= wfParams.surfaceW || gid.y >= wfParams.surfaceH) { return; } + + let pixelf = vec2(f32(gid.x), f32(gid.y)); + let res = vec2(f32(wfParams.surfaceW), f32(wfParams.surfaceH)); + let uv = (pixelf + vec2(0.5)) / res; + let ndc = uv * 2.0 - vec2(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(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 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 dim : Dim; +@group(1) @binding(0) var src : texture_2d; +@group(1) @binding(1) var dst : texture_storage_2d; + +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) gid: vec3) { + if (gid.x >= dim.w || gid.y >= dim.h) { return; } + let coord = vec2(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(0.2126, 0.7152, 0.0722)); + let keep = max(0.0, lum - 1.0) / max(lum, 1e-4); + textureStore(dst, coord, vec4(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).