Window::Render() rebuilt cmdBufInfo, the subresource range, both
VkImageMemoryBarrier structs, and presentInfo on the stack every frame
even though only the barriers' image (and the acquire barrier's
oldLayout) ever change. Hoist them to members initialised once in the
constructor; Render() now patches only the two varying fields.
presentInfo.pImageIndices/pSwapchains point at the currentBuffer and
swapChain members, so they track value changes (including swapchain
recreation) automatically. The render-complete wait semaphore is fixed
for the window's lifetime, so the per-frame `renderComplete != NULL`
check — set once, never cleared — was dead and is removed; the
semaphore is wired into presentInfo in the constructor instead.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Add zero-copy procedural BLAS overloads that take the AABB build input
straight from an existing device buffer instead of memcpy-ing a host span
through Mesh::aabbBuffer — the input-source companion to #36's refit work.
A GPU compute producer (e.g. a GPU-resident particle system) can now feed a
moving procedural BLAS that builds/refits each frame with no host round-trip.
Vulkan: new BuildProcedural(VkDeviceAddress, count, ...) and
RefitProcedural(VkDeviceAddress, count, ...) feed the device address directly
into VkAccelerationStructureGeometryAabbsDataKHR; the host-span path is
refactored to share RecordProceduralBuildFromAddress and is otherwise
unchanged.
WebGPU: device-buffer BuildProcedural/RefitProcedural copy the boxes GPU->GPU
into the mesh heap (new wgpuRegisterMeshBLASDeviceAabbs / wgpuRefitMeshBLAS-
DeviceAabbs bridge fns) and wrap them in a single root leaf bounded by a
caller-supplied worldBounds — no wasm round-trip, blasAddr stable across refit.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Mesh::Build / BuildProcedural now take an RTBuildOptions { preference,
allowUpdate }. `preference` maps to PREFER_FAST_TRACE (default) or
PREFER_FAST_BUILD; `allowUpdate` sets ALLOW_UPDATE so the BLAS can be
refit later.
Adds Mesh::Refit / RefitProcedural: when the original build opted into
allowUpdate and the topology is unchanged, they record an in-place
UPDATE-mode build (src == dst) — much cheaper than a rebuild, and the
AS handle + blasAddr are preserved so TLAS instances stay valid. They
fall back to a full rebuild otherwise. The shared build tail also now
destroys a stale AS handle on re-Build (previously leaked) and guards
the in-place update with an AS read→write barrier.
The portable RTBuildPreference/RTBuildOptions types and Refit methods
also exist on the WebGPU backend for API symmetry; the software BVH has
no hardware AS, so the preference is a no-op and a "refit" rebuilds the
BVH (re-registering the handle).
Also adds Device::validationErrorCount, bumped by the debug-messenger
callback on ERROR-severity messages, so tests can assert a Vulkan
operation produced no validation errors.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Adds the native counterpart of the WebGPU Mesh::BuildProcedural: uploads
the boxes as a 24-byte-stride VkAabbPositionsKHR-compatible device buffer
(AS-build-input read-only + device address) and builds the BLAS from one
VK_GEOMETRY_TYPE_AABBS_KHR geometry. opaque maps to
VK_GEOMETRY_OPAQUE_BIT_KHR, default-false semantics matching the WebGPU
path so any-hit shaders run. The BLAS sizing/create/build tail is factored
into a geometry-type-agnostic helper shared with the triangle path; no
TLAS/instance changes needed.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Sway delivers pointer events in surprising orders around input-device
hot-plug — observed: a second wl_pointer.leave after focus was already
cleared, which crashed in PointerListenerHandleLeave (null deref), and
motion/button would crash the same way if an enter never matched one of
our surfaces. Early-return instead. Found while end-to-end-testing the
scroll chain for #32 with a zwlr_virtual_pointer_v1 injector.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Window::onMouseScroll now fires on every backend and speaks one unit:
signed whole detents packed in the uint32 payload, +1 = wheel down
(DOM deltaY sign).
- Wayland: implement the previously-stubbed PointerListenerHandleAxis —
vertical axis only, wl_fixed_to_double / 15 (libinput's units-per-
detent), sub-detent remainder accumulated and reset on pointer leave.
- Win32: add the missing WM_MOUSEWHEEL case — -GET_WHEEL_DELTA_WPARAM /
WHEEL_DELTA with a remainder accumulator for free-spinning wheels.
- DOM: dom-env.js normalizes WheelEvent.deltaY per deltaMode (100px /
3 lines / 1 page per detent) with the same remainder scheme, so wasm
delivers the identical unit instead of raw browser-specific deltas.
- Document the contract on Window::onMouseScroll.
- tests/MouseScroll: compositor-free regression test driving the
Wayland axis handler directly.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
wfAccum becomes array<atomic<u32>> (4 slots/pixel, f32 bit patterns — same
16 B/pixel footprint) and rtAccumulate CASes each channel, so N rays for one
pixel may resolve in the same SHADE pass without the read-modify-write
racing. GENERATE clears with atomicStore (bitcast(0.0) == 0u); RESOLVE
atomicLoads + bitcasts the vec4 it hands runResolve.
Capacity half: RTPass::raysPerPixel (default 1) scales the wavefront
ray/hit/payload buffers to raysPerPixel·W·H rays per bounce so the
per-light emits actually fit instead of being dropped by rtEmitRay's
capacity guard. The accumulator stays per-pixel.
No flag-gating: uncontended the CAS succeeds first try — RTStress SHADE
stays at its documented ~1.0 ms and master-vs-branch renders are
byte-identical, so single-ray consumers pay nothing.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Adds the minimal WebGPU-backend surface an HDR bloom/post chain needs:
1. wgpuCreateStorageImage2D + StorageImage2D<>: runtime rgba16float (or
other format) textures with STORAGE_BINDING|TEXTURE_BINDING usage, no
CompressedTextureAsset upload — GPU-produced HDR targets.
2. UICustomBindingKind::StorageTexture (kind 5): write-only storage-texture
binding carrying a texel format (the freed UICustomBinding::format byte),
wired through all three binding paths (UI custom / RT / plain compute).
Float textures already sample via SampledTexture (sampleType float).
3. PipelineRTWebGPU::Init(..., hdrOutputFormat): RESOLVE writes the linear
accumulator into a user rgba16float texture (RTPass::outTexHandle)
instead of the rgba8unorm canvas, leaving the composite→swapchain pass to
the app. Default RGBA8Unorm keeps the canvas path byte-identical.
Existing RT examples updated for the _pad→format field rename; the default
paths are unchanged (verified RTStress builds + renders).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Two gaps in the Vulkan RT path that fault the device on the NVIDIA
proprietary driver with a non-trivial pipeline (simple VulkanTriangle
never hit them):
1. maxPipelineRayRecursionDepth was hardcoded to 1, so any closest-hit
shader that traces a secondary ray (shadow ray — a very common
pattern) recursed past the pipeline limit (UB → device fault).
PipelineRTVulkan::Init now takes a maxRecursionDepth parameter
(default 1, clamped to the device's maxRayRecursionDepth).
2. The NVIDIA descriptor-heap AS-read workaround rewrites every shader
that reads an accelerationStructureEXT from the heap — including
compute shaders — to read the TLAS device address from a push
constant, but only RTPass pushed that address. A compute shader that
ray-queries the TLAS (rayQueryEXT) therefore ran against an unwritten
push slot → garbage AS handle → VK_ERROR_DEVICE_LOST.
WorkaroundNvidiaAS::Patch now returns a per-shader PatchResult
{patched, tlasPushOffset} instead of writing the clobber-prone global
Device::workaroundTlasPushOffset (removed). VulkanShader stores it;
ShaderBindingTableVulkan/PipelineRTVulkan carry it for RTPass, and
ComputeShader tracks its own offset and pushes the caller-supplied
TLAS address in Dispatch (new defaulted tlasAddress parameter),
mirroring RTPass::Record.
The PushConstantRewrite regression test now asserts Patch's returned
patched/offset and adds two ray-querying compute-shader cases, proving
the rewrite is stage-agnostic and the per-shader offset is correct.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The GPU-AV enable list was removed to dodge a crash in SDK 1.4.341,
whose GPU-AV null-deref'd on descriptor_heap pipelines
(VK_PIPELINE_CREATE_2_DESCRIPTOR_HEAP_BIT_EXT, layout = VK_NULL_HANDLE)
in PipelineSubState::GetPipelineLayoutUnion:
https://github.com/KhronosGroup/Vulkan-ValidationLayers/issues/12103
That was fixed in the next SDK release. The validation layer is now
1.4.350 (> 1.4.341), so restore VK_VALIDATION_FEATURE_ENABLE_GPU_ASSISTED_EXT
in the VkValidationFeaturesEXT enable list.
Verified by running the HelloUI example (which draws through the
descriptor_heap compute pipelines) with the layer active: it renders the
full UI for the entire run with GPU-AV reporting "Both GPU Assisted
Validation and Normal Core Check Validation are enabled" and no
descriptor-heap null-deref or VUID errors.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Reading an acceleration structure through VK_EXT_descriptor_heap aborts
with VK_ERROR_DEVICE_LOST on NVIDIA 610.43.02 — a brand-new-extension
driver fault isolated in #7 (engine setup is correct and validation-clean;
images/buffers through the same heap work, and both traceRayEXT and inline
rayQuery fault identically on the AS read).
An acceleration structure can equally be reached by its device address via
OpConvertUToAccelerationStructureKHR, which reads no descriptor and so never
touches the faulting heap path. glslang has no GLSL spelling for that
conversion, so VulkanShader rewrites the compiled SPIR-V at module-load
time: every `OpLoad %accelStruct <heap-ptr>` becomes a load of the TLAS
device address from a synthesized push-constant block followed by the
convert. RTPass pushes the active frame's TLAS address into that push
constant. User GLSL and example code are unchanged; acceleration structures
still bind into the heap normally.
The workaround is gated on Device::workaroundDescriptorHeapAS (true only on
the NVIDIA proprietary driver) and confined to one fenced block in
Crafter.Graphics-ShaderVulkan.cppm plus the RTPass push and the shaderInt64
feature toggle — delete those once a fixed NVIDIA driver ships and the heap
AS read becomes the direct path again.
Verified: VulkanTriangle ray-traces correctly on native NVIDIA (RTX 4090),
validation-layer-clean, no device loss. The SPIR-V rewrite was independently
validated with spirv-val on both the VulkanTriangle and Sponza raygen
modules.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
PipelineRTWebGPU emits a runIntersection mega-switch and the
RT_HAS_ANYHIT / RT_HAS_INTERSECTION consts (+ the @CRAFTER_RT_TRACE_USER
marker) that gate the library's new TRACE-stage user callbacks, so an
opaque triangle-only scene still const-folds them away. Mesh-WebGPU
builds a SAH BVH2 over AABB primitives and uploads them in primitive
order for the intersection shader to fetch.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
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 <noreply@anthropic.com>
Two Vulkan validation errors fired on startup of every native (Vulkan)
example, reported in #5:
1. vkCreateDevice enabledLayerCount != 0. Device layers are deprecated
and ignored since Vulkan 1.0; passing them is a spec violation
(VUID-VkDeviceCreateInfo-enabledLayerCount-12384). The device-layer
enumeration/match block in Device::Initialize is removed and
enabledLayerCount is pinned to 0 — layers are enabled at the instance
only.
2. vkQueueSubmit layout transition on a presentable image that "has not
been acquired". StartInit() and RecreateSwapchainAndImages() eagerly
transitioned every swapchain image UNDEFINED -> PRESENT_SRC_KHR before
any vkAcquireNextImageKHR, which the spec forbids (a presentable image
may only be touched after acquire). Those pre-transitions are removed.
Each image's first layout transition now happens lazily in Render(),
after acquire, from UNDEFINED; subsequent frames transition from
PRESENT_SRC_KHR. A per-image `imageInitialised` flag (reset in
CreateSwapchain) selects the correct oldLayout.
Verified under sway (headless, GPU renderer) + VK_LAYER_KHRONOS_validation:
the original code reproduces both errors on HelloUI; the fixed build emits
zero validation messages across initial render and swapchain recreation.
Resolves#5
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
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>
