4e42d663a6
Replace the megakernel @compute entry with five wavefront kernels sharing one module, connected by GPU ray/hit/payload buffers and a GPU-driven indirect bounce loop: GENERATE -> (PREP -> TRACE -> SHADE) x maxDepth -> RESOLVE - TRACE contains zero user code (pure _rtwTraverseTlas/Blas, opaque-only). - PREP publishes dispatchWorkgroupsIndirect args from the live ray count; the indirect-args buffer lives in its own bind group so it is never bound read-write in the same dispatch that consumes it as INDIRECT. - New emit/accumulate API: rtEmitPrimaryRay / rtEmitRay / rtAccumulate, plus an optional user Resolve stage (tonemap hook; identity by default). - Per-pass WfParams via a dynamic-offset uniform ring (curIsA/bounce vary between passes within one submit). - Payload-typed wfPayload binding emitted in the codegen region after the user's struct Payload; payload travels with each ray (2*W*H slots). - Request maxBufferSize / maxStorageBufferBindingSize / maxComputeWorkgroups PerDimension so the W*H-sized work buffers fit past the 128MB baseline. VulkanTriangle ported to the new API and renders bit-identical to the megakernel baseline at maxDepth=1. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
34 lines
1.1 KiB
WebGPU Shading Language
34 lines
1.1 KiB
WebGPU Shading Language
// WebGPU wavefront raygen. Runs in GENERATE: compute the pinhole camera
|
|
// ray and emit it as the pixel's primary ray. Shading happens later in
|
|
// SHADE (closesthit/miss). The Payload type is declared in closesthit.wgsl.
|
|
|
|
fn raygen_main(gid: vec3<u32>) {
|
|
if (gid.x >= wfParams.surfaceW || gid.y >= wfParams.surfaceH) { return; }
|
|
|
|
let pixel = vec2<f32>(f32(gid.x), f32(gid.y));
|
|
let resolution = vec2<f32>(f32(wfParams.surfaceW), f32(wfParams.surfaceH));
|
|
let uv = (pixel + vec2<f32>(0.5)) / resolution;
|
|
let ndc = uv * 2.0 - vec2<f32>(1.0);
|
|
|
|
let origin = vec3<f32>(0.0, 0.0, -300.0);
|
|
let aspect = resolution.x / resolution.y;
|
|
let fov = 60.0 * 3.14159265 / 180.0;
|
|
let tanHalfFov = tan(fov * 0.5);
|
|
|
|
let direction = normalize(vec3<f32>(
|
|
ndc.x * aspect * tanHalfFov,
|
|
-ndc.y * tanHalfFov,
|
|
1.0,
|
|
));
|
|
|
|
var payload: Payload;
|
|
payload.color = vec3<f32>(0.0);
|
|
|
|
rtEmitPrimaryRay(
|
|
origin, 0.001,
|
|
direction, 10000.0,
|
|
0u, // ray flags
|
|
0xFFu, // cull mask
|
|
0u, 0u, // sbtRecordOffset, missIndex
|
|
payload);
|
|
}
|