From 1e749818ef6c4e3878beae4f3773d519bb2e3b93 Mon Sep 17 00:00:00 2001 From: catbot Date: Mon, 1 Jun 2026 11:09:15 +0000 Subject: [PATCH] fix(webgpu): reshape wavefront TRACE/SHADE to 2-D to survive >4.19M rays MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A 1-D indirect dispatch of ceil(W*H/64) workgroups for the wavefront TRACE/SHADE stages overflows maxComputeWorkgroupsPerDimension (65535 on Dawn/Firefox) once the surface exceeds ~4.19M rays (~2560x1640). Per the WebGPU spec such a dispatch is silently dropped — no validation error — so at 4K the world is never traced and the accumulator stays black while non-RT passes survive. _wfPrep now spreads the workgroups across a 2-D grid (x clamped to 65535, y = ceil(wg/65535)), and the wfTrace/wfShade entry points rebuild the linear ray index from (global_invocation_id, num_workgroups). The existing `i >= _wfCurCount()` guard absorbs the grid overshoot. GENERATE/RESOLVE already use a 2-D tile dispatch and are unchanged. Verified in Firefox/WebGPU with RTStress at a 3449x1739 surface (5.99M rays, 93716 workgroups — well over the 65535 cap): renders the full cube grid where master shows a black screen. Resolves #11 Co-Authored-By: Claude Opus 4.8 --- additional/dom-webgpu.js | 15 ++++++++++++--- .../Crafter.Graphics-PipelineRTWebGPU.cpp | 4 ++-- 2 files changed, 14 insertions(+), 5 deletions(-) diff --git a/additional/dom-webgpu.js b/additional/dom-webgpu.js index 625a552..a15142f 100644 --- a/additional/dom-webgpu.js +++ b/additional/dom-webgpu.js @@ -1834,9 +1834,18 @@ fn _wfReadRay(i: u32) -> WfRay { // PREP — publish indirect args for the upcoming TRACE/SHADE; zero the next // buffer's emit counter. fn _wfPrep() { - let n = _wfCurCount(); - wfIndirect[0] = (n + 63u) / 64u; - wfIndirect[1] = 1u; + let n = _wfCurCount(); + let wg = (n + 63u) / 64u; + // maxComputeWorkgroupsPerDimension is 65535 on Dawn/Firefox; a 1-D + // dispatch of ceil(W*H/64) overflows it past ~4.19M rays (~2560x1640) + // and WebGPU silently drops the indirect dispatch -> black screen. Spread + // across a 2-D grid; wfTrace/wfShade rebuild the linear index from + // num_workgroups. + let MAXDIM = 65535u; + let gx = min(wg, MAXDIM); + let gy = (wg + MAXDIM - 1u) / MAXDIM; // = 1 when wg <= MAXDIM + wfIndirect[0] = gx; + wfIndirect[1] = gy; wfIndirect[2] = 1u; if (wfParams.curIsA == 1u) { atomicStore(&wfCounters[1], 0u); } else { atomicStore(&wfCounters[0], 0u); } diff --git a/implementations/Crafter.Graphics-PipelineRTWebGPU.cpp b/implementations/Crafter.Graphics-PipelineRTWebGPU.cpp index 8dd949d..373249a 100644 --- a/implementations/Crafter.Graphics-PipelineRTWebGPU.cpp +++ b/implementations/Crafter.Graphics-PipelineRTWebGPU.cpp @@ -231,13 +231,13 @@ void PipelineRTWebGPU::Init(WebGPUCommandEncoderRef /*cmd*/, // TRACE — zero user code: pure traversal + intersection. One thread // per live ray; writes a HitResult into wfHits[i]. wgsl += "\n@compute @workgroup_size(64)\n"; - wgsl += "fn wfTrace(@builtin(global_invocation_id) gid: vec3) { _wfTrace(gid.x); }\n"; + wgsl += "fn wfTrace(@builtin(global_invocation_id) gid: vec3, @builtin(num_workgroups) nwg: vec3) { _wfTrace(gid.y * nwg.x * 64u + gid.x); }\n"; // SHADE — one thread per live ray; loads the ray + its hit + payload, // dispatches to runMiss / runClosestHit, which may rtAccumulate and // rtEmitRay continuation/shadow rays into the next buffer. wgsl += "\n@compute @workgroup_size(64)\n"; - wgsl += "fn wfShade(@builtin(global_invocation_id) gid: vec3) { _wfShade(gid.x); }\n"; + wgsl += "fn wfShade(@builtin(global_invocation_id) gid: vec3, @builtin(num_workgroups) nwg: vec3) { _wfShade(gid.y * nwg.x * 64u + gid.x); }\n"; // RESOLVE — one thread per pixel; runs the user resolve (or identity) // over the linear accumulator and stores to the output image. -- 2.54.0