bbe1b21c22
RT→rgba16float→threshold→blur→composite, end-to-end proof of the three primitives. threshold/blur run from onBeforeUpdate (one submit each) so the storage-write→sampled-read dependency gets WebGPU's per-submit barrier (there is none between dispatches within a compute pass); the composite is a UI custom shader so it owns the canvas ping-pong. Documents the Vulkan symmetry (gap 4): the native present path records passes generically and barriers between them, so the same chain is wireable today. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
53 lines
2.3 KiB
WebGPU Shading Language
53 lines
2.3 KiB
WebGPU Shading Language
// 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));
|
|
}
|