From f14441942a8c9cef882df71e906de4928c637f4f Mon Sep 17 00:00:00 2001 From: catbot Date: Tue, 16 Jun 2026 14:12:04 +0000 Subject: [PATCH 1/3] feat(rt): BuildProcedural/RefitProcedural from a device AABB buffer (#37) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- additional/dom-webgpu.js | 124 +++++++++++++++++- .../Crafter.Graphics-Mesh-WebGPU.cpp | 38 ++++++ implementations/Crafter.Graphics-Mesh.cpp | 79 ++++++++--- interfaces/Crafter.Graphics-Mesh.cppm | 56 ++++++++ interfaces/Crafter.Graphics-WebGPU.cppm | 30 +++++ 5 files changed, 306 insertions(+), 21 deletions(-) diff --git a/additional/dom-webgpu.js b/additional/dom-webgpu.js index 9d76441..c35456b 100644 --- a/additional/dom-webgpu.js +++ b/additional/dom-webgpu.js @@ -51,13 +51,14 @@ function stub(name) { "wgpuFrameBegin", "wgpuFrameEnd", "wgpuDispatchQuads", "wgpuDispatchCircles", "wgpuDispatchImages", "wgpuDispatchText", "wgpuLoadCustomShader", "wgpuDispatchCustom", - "wgpuRegisterMeshBLAS", "wgpuLoadRTPipeline", "wgpuDispatchRT", "wgpuBuildTLAS", + "wgpuRegisterMeshBLAS", "wgpuRegisterMeshBLASDeviceAabbs", "wgpuRefitMeshBLASDeviceAabbs", + "wgpuLoadRTPipeline", "wgpuDispatchRT", "wgpuBuildTLAS", "wgpuLoadComputePipeline", "wgpuDispatchCompute", ]) { // Read-write ints don't need a stub-throw; return 0 for the size queries. e[n] = n.endsWith("Width") || n.endsWith("Height") ? () => 0 - : (n === "wgpuRegisterMeshBLAS" ? () => 0 : stub(n)); + : ((n === "wgpuRegisterMeshBLAS" || n === "wgpuRegisterMeshBLASDeviceAabbs") ? () => 0 : stub(n)); } } @@ -2761,6 +2762,12 @@ const rtState = { // compute shaders at dispatch time. currentTlas: 0, currentTlasInstanceCount: 0, + + // handle → { vCursor, bvhNode, primCount } for device-buffer procedural + // BLAS (wgpuRegisterMeshBLASDeviceAabbs). Lets the per-frame refit copy + // fresh boxes into the same vertices-heap region and rewrite the root + // leaf bounds without re-registering the mesh. + deviceAabbRegions: new Map(), }; function rtInit() { @@ -2933,6 +2940,119 @@ env.wgpuRegisterMeshBLAS = (minX, minY, minZ, maxX, maxY, maxZ, return handle; }; +// Build a single-leaf BVH node (32 bytes) spanning all `count` prims. +function _rtWriteRootLeaf(bvhCursorBytes, count, minX, minY, minZ, maxX, maxY, maxZ) { + const node = new ArrayBuffer(32); + const nf = new Float32Array(node); + const nu = new Uint32Array(node); + nf[0] = minX; nf[1] = minY; nf[2] = minZ; + nu[3] = 0; // firstChildOrPrim: prim base (rel to primRemapOffset) + nf[4] = maxX; nf[5] = maxY; nf[6] = maxZ; + nu[7] = count >>> 0; // primCount > 0 → leaf + queue.writeBuffer(rtState.bvhHeap.gpu, bvhCursorBytes, node); +} + +// Zero-copy procedural BLAS: the boxes already live in a device GPUBuffer +// (a compute pass wrote them). Copy them GPU→GPU into the vertices heap and +// register a single-root-leaf BLAS — no wasm round-trip, no SAH build (there +// is no host copy to build over). See wgpuRegisterMeshBLASDeviceAabbs decl in +// Crafter.Graphics-WebGPU.cppm. +env.wgpuRegisterMeshBLASDeviceAabbs = (aabbBufHandle, count, + minX, minY, minZ, maxX, maxY, maxZ, + opaqueFlag) => { + if (!rtState.vertHeap) rtInit(); + const src = buffers.get(aabbBufHandle); + if (!src) { + console.error("[crafter-wgpu] wgpuRegisterMeshBLASDeviceAabbs: unknown buffer handle"); + return 0; + } + console.log(`[crafter-wgpu] device-AABB BLAS: ${count} aabbs, bbox=(${minX.toFixed(1)}..${maxX.toFixed(1)}, ${minY.toFixed(1)}..${maxY.toFixed(1)}, ${minZ.toFixed(1)}..${maxZ.toFixed(1)}), opaque=${opaqueFlag}`); + + const vBytes = count * 24; // 2 vec3 per box, matching RTAabb/VkAabbPositionsKHR + const nBytes = 32; // one root leaf + const rBytes = count * 4; // identity primRemap + + rtHeapEnsure(rtState.vertHeap, vBytes); + rtHeapEnsure(rtState.bvhHeap, nBytes); + rtHeapEnsure(rtState.primRemapHeap, rBytes); + + const vCursor = rtState.vertHeap.cursor; + const nCursor = rtState.bvhHeap.cursor; + const rCursor = rtState.primRemapHeap.cursor; + const vOff = vCursor / 12; // in vec3 units + const nOff = nCursor / 32; // in BVHNode units + const rOff = rCursor / 4; // in u32 units + + // GPU→GPU copy of the boxes — the data never touches wasm memory. + { + const enc = device.createCommandEncoder(); + enc.copyBufferToBuffer(src, 0, rtState.vertHeap.gpu, vCursor, vBytes); + queue.submit([enc.finish()]); + } + + // Identity primRemap [0..count) — leaf slot i → box i in the heap. + const remap = new Uint32Array(count); + for (let i = 0; i < count; i++) remap[i] = i; + queue.writeBuffer(rtState.primRemapHeap.gpu, rCursor, remap); + + _rtWriteRootLeaf(nCursor, count, minX, minY, minZ, maxX, maxY, maxZ); + + rtState.vertHeap.cursor += vBytes; + rtState.bvhHeap.cursor += nBytes; + rtState.primRemapHeap.cursor += rBytes; + + const handle = rtState.nextMeshHandle++; + rtMeshRecordsEnsure(handle + 1); + + const rec = new ArrayBuffer(64); + const f32 = new Float32Array(rec); + const u32 = new Uint32Array(rec); + f32[0] = minX; f32[1] = minY; f32[2] = minZ; + u32[3] = vOff; + f32[4] = maxX; f32[5] = maxY; f32[6] = maxZ; + u32[7] = 0; // indexOffset (unused for AABBs) + u32[8] = nOff; + u32[9] = rOff; + u32[10] = count >>> 0; // primitive count + u32[11] = 0; // attribsOffset (none) + u32[12] = 1; // geomType = AABBs + u32[13] = opaqueFlag ? 1 : 0; + u32[14] = 0; + u32[15] = 0; + queue.writeBuffer(rtState.meshRecordsBuffer, handle * 64, rec); + + rtState.deviceAabbRegions.set(handle, { vCursor, nCursor, primCount: count }); + return handle; +}; + +// Zero-copy procedural refit: re-copy the boxes into the existing heap +// region and refresh the root leaf bounds. Handle (and thus blasAddr) is +// preserved. Count must match the original register. +env.wgpuRefitMeshBLASDeviceAabbs = (meshHandle, aabbBufHandle, count, + minX, minY, minZ, maxX, maxY, maxZ) => { + const region = rtState.deviceAabbRegions.get(meshHandle); + const src = buffers.get(aabbBufHandle); + if (!region || !src) { + console.error("[crafter-wgpu] wgpuRefitMeshBLASDeviceAabbs: unknown mesh/buffer handle"); + return; + } + if (count !== region.primCount) { + console.error(`[crafter-wgpu] wgpuRefitMeshBLASDeviceAabbs: count ${count} != built ${region.primCount}`); + return; + } + const enc = device.createCommandEncoder(); + enc.copyBufferToBuffer(src, 0, rtState.vertHeap.gpu, region.vCursor, count * 24); + queue.submit([enc.finish()]); + // Refresh the root leaf bounds so traversal's node-level cull stays valid. + _rtWriteRootLeaf(region.nCursor, count, minX, minY, minZ, maxX, maxY, maxZ); + // Refresh the MeshRecord root AABB (used by TLAS build for the world box). + const hdr = new Float32Array(8); + hdr[0] = minX; hdr[1] = minY; hdr[2] = minZ; + new Uint32Array(hdr.buffer)[3] = region.vCursor / 12; + hdr[4] = maxX; hdr[5] = maxY; hdr[6] = maxZ; + queue.writeBuffer(rtState.meshRecordsBuffer, meshHandle * 64, hdr.buffer, 0, 32); +}; + env.wgpuBuildTLAS = (instanceBufHandle, instanceCount, tlasOutBufHandle, entryOrderHandle, mortonHandle, binsHandle, bvhNodesHandle, sortABufHandle, sortBBufHandle) => { diff --git a/implementations/Crafter.Graphics-Mesh-WebGPU.cpp b/implementations/Crafter.Graphics-Mesh-WebGPU.cpp index 5d74198..e85317c 100644 --- a/implementations/Crafter.Graphics-Mesh-WebGPU.cpp +++ b/implementations/Crafter.Graphics-Mesh-WebGPU.cpp @@ -342,6 +342,28 @@ void Mesh::BuildProcedural(std::span aabbs, /*primCount*/ static_cast(count)); } +void Mesh::BuildProcedural(WebGPUBufferRef aabbBuffer, + std::uint32_t count, + RTAabb worldBounds, + bool opaque_, + WebGPUCommandEncoderRef /*cmd*/, + RTBuildOptions /*options*/) { + // Zero-copy: the boxes are already on the GPU (a compute pass wrote + // them). The bridge copies them GPU→GPU into the vertices heap and + // registers a single-root-leaf BLAS bounded by worldBounds — no SAH + // build (there is no host copy of the boxes to build over) and no host + // round-trip. Traversal linearly intersects all `count` boxes. + opaque = opaque_; + triangleCount = 0; // not a triangle mesh + vertexCount = count * 2; // 2 "vertices" (min,max) per box + + blasAddr = WebGPU::wgpuRegisterMeshBLASDeviceAabbs( + aabbBuffer, static_cast(count), + worldBounds.min[0], worldBounds.min[1], worldBounds.min[2], + worldBounds.max[0], worldBounds.max[1], worldBounds.max[2], + opaque ? 1 : 0); +} + void Mesh::Refit(std::span> vertices, std::span indices, WebGPUCommandEncoderRef cmd) { @@ -359,3 +381,19 @@ void Mesh::RefitProcedural(std::span aabbs, WebGPUCommandEncoderRef cmd) { BuildProcedural(aabbs, opaque, cmd); } + +void Mesh::RefitProcedural(WebGPUBufferRef aabbBuffer, + std::uint32_t count, + RTAabb worldBounds, + WebGPUCommandEncoderRef /*cmd*/) { + // Re-copy the GPU boxes into the existing heap region and refresh the + // root bounds — keeps blasAddr stable (no re-register), so TLAS + // instances referencing this mesh stay valid across the per-frame + // update. No host copy. + vertexCount = count * 2; + WebGPU::wgpuRefitMeshBLASDeviceAabbs( + static_cast(blasAddr), aabbBuffer, + static_cast(count), + worldBounds.min[0], worldBounds.min[1], worldBounds.min[2], + worldBounds.max[0], worldBounds.max[1], worldBounds.max[2]); +} diff --git a/implementations/Crafter.Graphics-Mesh.cpp b/implementations/Crafter.Graphics-Mesh.cpp index 2dbff16..203e028 100644 --- a/implementations/Crafter.Graphics-Mesh.cpp +++ b/implementations/Crafter.Graphics-Mesh.cpp @@ -267,26 +267,18 @@ void Mesh::Build(const CompressedMeshAsset& asset, VkCommandBuffer cmd, RTBuildO } namespace { - // Re-upload AABB build input and record an AABB BLAS build (fresh or - // in-place refit). Shared by BuildProcedural and RefitProcedural; the - // geometry's opaque bit is read from self.opaque so an UPDATE keeps the - // exact geometry description of the original build. - void RecordProceduralBuild(Mesh& self, std::span aabbs, VkBuildAccelerationStructureFlagsKHR flags, bool update, VkCommandBuffer cmd) { - // 24-byte-stride VkAabbPositionsKHR-compatible build input - // (static_assert'd in the interface). Same usage set as the triangle - // inputs: AS-build read-only + device address. A refit reuses the - // existing same-sized buffer (count is unchanged), so the device - // address — and the geometry it feeds — stays stable. - if (!update) { - self.aabbBuffer.Resize(kVertexUsageBase, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT, static_cast(aabbs.size())); - } - std::memcpy(self.aabbBuffer.value, aabbs.data(), aabbs.size() * sizeof(RTAabb)); - self.aabbBuffer.FlushDevice(cmd, VK_ACCESS_MEMORY_READ_BIT, VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR); - + // Record an AABB BLAS build (fresh or in-place refit) from a device + // address that already holds the VkAabbPositionsKHR-compatible boxes. + // Geometry-source agnostic: the host-upload path (RecordProceduralBuild) + // points this at self.aabbBuffer, while the device-buffer overloads + // point it straight at the producer's buffer — no copy involved either + // way. The geometry's opaque bit is read from self.opaque so an UPDATE + // keeps the exact geometry description of the original build. + void RecordProceduralBuildFromAddress(Mesh& self, VkDeviceAddress aabbAddress, std::uint32_t count, std::uint32_t stride, VkBuildAccelerationStructureFlagsKHR flags, bool update, VkCommandBuffer cmd) { VkAccelerationStructureGeometryAabbsDataKHR aabbsData { .sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_AABBS_DATA_KHR, - .data = { .deviceAddress = self.aabbBuffer.address }, - .stride = sizeof(RTAabb) + .data = { .deviceAddress = aabbAddress }, + .stride = stride }; VkAccelerationStructureGeometryDataKHR geometryData; geometryData.aabbs = aabbsData; @@ -299,7 +291,25 @@ namespace { .flags = self.opaque ? static_cast(VK_GEOMETRY_OPAQUE_BIT_KHR) : VkGeometryFlagsKHR{} }; - RecordBLASBuildFromGeometry(self, blasGeometry, static_cast(aabbs.size()), flags, update, cmd); + RecordBLASBuildFromGeometry(self, blasGeometry, count, flags, update, cmd); + } + + // Re-upload AABB build input from host memory and record an AABB BLAS + // build (fresh or in-place refit). Shared by the host-span BuildProcedural + // and RefitProcedural. + void RecordProceduralBuild(Mesh& self, std::span aabbs, VkBuildAccelerationStructureFlagsKHR flags, bool update, VkCommandBuffer cmd) { + // 24-byte-stride VkAabbPositionsKHR-compatible build input + // (static_assert'd in the interface). Same usage set as the triangle + // inputs: AS-build read-only + device address. A refit reuses the + // existing same-sized buffer (count is unchanged), so the device + // address — and the geometry it feeds — stays stable. + if (!update) { + self.aabbBuffer.Resize(kVertexUsageBase, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT, static_cast(aabbs.size())); + } + std::memcpy(self.aabbBuffer.value, aabbs.data(), aabbs.size() * sizeof(RTAabb)); + self.aabbBuffer.FlushDevice(cmd, VK_ACCESS_MEMORY_READ_BIT, VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR); + + RecordProceduralBuildFromAddress(self, self.aabbBuffer.address, static_cast(aabbs.size()), sizeof(RTAabb), flags, update, cmd); } } @@ -311,6 +321,20 @@ void Mesh::BuildProcedural(std::span aabbs, bool opaque, VkCommand RecordProceduralBuild(*this, aabbs, BlasFlags(options), /*update*/ false, cmd); } +void Mesh::BuildProcedural(VkDeviceAddress aabbAddress, std::uint32_t count, bool opaque, VkCommandBuffer cmd, RTBuildOptions options, std::uint32_t stride) { + // Zero-copy procedural build: the AABBs already live in `aabbAddress` + // (a device buffer a GPU compute pass wrote). Nothing touches + // self.aabbBuffer — the build input is the producer's buffer directly. + // The caller owns that buffer's lifetime + usage flags + // (ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY + SHADER_DEVICE_ADDRESS) + // and must barrier the producing writes before this build reads them. + this->opaque = opaque; + allowUpdate = options.allowUpdate; + builtInputCount = count; + + RecordProceduralBuildFromAddress(*this, aabbAddress, count, stride, BlasFlags(options), /*update*/ false, cmd); +} + void Mesh::Refit(std::span> verticies, std::span indicies, VkCommandBuffer cmd) { // A hardware in-place UPDATE is only valid when the original build asked // for it, an AS already exists, and the topology is unchanged. Otherwise @@ -356,3 +380,20 @@ void Mesh::RefitProcedural(std::span aabbs, VkCommandBuffer cmd) { RecordProceduralBuild(*this, aabbs, buildFlags, update, cmd); } +void Mesh::RefitProcedural(VkDeviceAddress aabbAddress, std::uint32_t count, VkCommandBuffer cmd, std::uint32_t stride) { + // Zero-copy refit: the GPU producer rewrote the same device buffer in + // place; re-record the build (UPDATE when allowed + count unchanged, + // else a full rebuild) reading straight from `aabbAddress`. As with the + // device-buffer BuildProcedural, nothing touches self.aabbBuffer and the + // caller is responsible for barriering the producing writes. + const bool sameTopology = + accelerationStructure != VK_NULL_HANDLE + && count == builtInputCount; + const bool update = allowUpdate && sameTopology; + + if (!update && count != builtInputCount) { + builtInputCount = count; + } + RecordProceduralBuildFromAddress(*this, aabbAddress, count, stride, buildFlags, update, cmd); +} + diff --git a/interfaces/Crafter.Graphics-Mesh.cppm b/interfaces/Crafter.Graphics-Mesh.cppm index 516d408..c04f819 100644 --- a/interfaces/Crafter.Graphics-Mesh.cppm +++ b/interfaces/Crafter.Graphics-Mesh.cppm @@ -125,6 +125,23 @@ export namespace Crafter { // exactly like a triangle BLAS. void BuildProcedural(std::span aabbs, bool opaque, VkCommandBuffer cmd, RTBuildOptions options = {}); + // Zero-copy procedural build: take the AABB build input straight from + // an existing device buffer (by device address + primitive count) + // rather than memcpy-ing a host span through aabbBuffer. The intended + // producer is a GPU compute pass that writes the boxes itself — e.g. + // a GPU-resident particle system whose per-frame AABBs never touch the + // CPU. The buffer must carry + // VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR + // and VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT, its boxes laid out as + // VkAabbPositionsKHR (`stride`, default sizeof(RTAabb) == 24), and the + // caller must pipeline-barrier the producing writes (e.g. a compute + // dispatch) before the build on `cmd` reads them. The buffer's + // lifetime is the caller's responsibility — it is never copied, so it + // must outlive the build submitted on `cmd`. Everything else (opaque + // bit, hit-group contract, instance wiring) matches the host-span + // BuildProcedural above. + void BuildProcedural(VkDeviceAddress aabbAddress, std::uint32_t count, bool opaque, VkCommandBuffer cmd, RTBuildOptions options = {}, std::uint32_t stride = sizeof(RTAabb)); + // Refit the triangle BLAS against new geometry of the *same // topology* (same vertex count and index count as the original // Build). When that Build set RTBuildOptions::allowUpdate this issues @@ -139,6 +156,16 @@ export namespace Crafter { void Refit(std::span> verticies, std::span indicies, VkCommandBuffer cmd); // Procedural analog of Refit: new object-space boxes, same count. void RefitProcedural(std::span aabbs, VkCommandBuffer cmd); + // Zero-copy procedural refit: the device-buffer counterpart of + // RefitProcedural above and the input-source pairing for the device + // BuildProcedural. The GPU producer rewrote the boxes in `aabbAddress` + // in place; this re-records the build straight from that address — + // an in-place hardware UPDATE when the original build set allowUpdate + // and `count` is unchanged (handle / blasAddr preserved), otherwise a + // full rebuild. Same buffer-usage / barrier / lifetime contract as the + // device BuildProcedural. Call per frame to track GPU-written boxes + // with zero host involvement. + void RefitProcedural(VkDeviceAddress aabbAddress, std::uint32_t count, VkCommandBuffer cmd, std::uint32_t stride = sizeof(RTAabb)); }; } #endif // !CRAFTER_GRAPHICS_WINDOW_DOM @@ -243,6 +270,25 @@ export namespace Crafter { WebGPUCommandEncoderRef cmd = 0, RTBuildOptions options = {}); + // Zero-copy procedural build: the boxes already live in an existing + // device buffer (`aabbBuffer`, a GPUBuffer a compute pass wrote), fed + // straight into the BLAS with no host copy — the WebGPU analog of the + // device-address Vulkan BuildProcedural. The software path has no + // hardware AS and builds its BVH on the host, but here there is no + // host copy to build over, so the device boxes are copied GPU→GPU into + // the mesh heap and wrapped in a single root leaf bounded by + // `worldBounds` (object-space min/max enclosing all `count` boxes — + // the producer supplies it since the CPU never sees the boxes). The + // traversal then linearly intersects every box. `opaque` is the + // geometry opaque bit (false lets any-hit run). `options` has no effect + // (no hardware AS to tune), kept for API symmetry. + void BuildProcedural(WebGPUBufferRef aabbBuffer, + std::uint32_t count, + RTAabb worldBounds, + bool opaque = false, + WebGPUCommandEncoderRef cmd = 0, + RTBuildOptions options = {}); + // Refit analogs of the native API. With no hardware AS to update, // these simply re-run the host BVH build over the new data, so they // accept the full geometry (not just moved positions). Provided so @@ -253,6 +299,16 @@ export namespace Crafter { WebGPUCommandEncoderRef cmd = 0); void RefitProcedural(std::span aabbs, WebGPUCommandEncoderRef cmd = 0); + // Zero-copy procedural refit: re-copy the boxes from the device buffer + // into this mesh's heap region (count unchanged) and update the root + // leaf bounds, keeping blasAddr stable — unlike the host-span + // RefitProcedural, which rebuilds and re-publishes a NEW handle. The + // device counterpart of the device-address Vulkan RefitProcedural; + // call per frame to track GPU-written boxes with no host copy. + void RefitProcedural(WebGPUBufferRef aabbBuffer, + std::uint32_t count, + RTAabb worldBounds, + WebGPUCommandEncoderRef cmd = 0); }; } #endif // CRAFTER_GRAPHICS_WINDOW_DOM diff --git a/interfaces/Crafter.Graphics-WebGPU.cppm b/interfaces/Crafter.Graphics-WebGPU.cppm index a7ea9d2..0cc81c9 100644 --- a/interfaces/Crafter.Graphics-WebGPU.cppm +++ b/interfaces/Crafter.Graphics-WebGPU.cppm @@ -195,6 +195,36 @@ namespace Crafter::WebGPU { const void* attribsPtr, std::int32_t attribsByteCount, std::int32_t geomType, std::int32_t opaqueFlag, std::int32_t primCount); + // Zero-copy procedural BLAS registration: the AABB build input already + // lives in an existing device buffer (`aabbBufferHandle`, a GPUBuffer a + // compute pass wrote) rather than a wasm host pointer. The bridge + // copyBufferToBuffer's `count` boxes (24 bytes each, [min,max] vec3 — the + // RTAabb / VkAabbPositionsKHR layout) straight into the vertices heap, + // never round-tripping them through wasm memory. With no host copy of the + // boxes there is nothing to run the SAH builder over, so the BLAS gets a + // single root leaf spanning all `count` primitives bounded by the + // caller-supplied object-space box (min/max) — the traversal then linearly + // intersects every box. `opaqueFlag` is the geometry opaque bit. Returns + // the mesh handle (0 on failure), same contract as wgpuRegisterMeshBLAS. + __attribute__((import_module("env"), import_name("wgpuRegisterMeshBLASDeviceAabbs"))) + extern "C" std::uint32_t wgpuRegisterMeshBLASDeviceAabbs( + WebGPUBufferRef aabbBufferHandle, std::int32_t count, + float minX, float minY, float minZ, + float maxX, float maxY, float maxZ, + std::int32_t opaqueFlag); + + // Zero-copy procedural BLAS refit: re-copy `count` boxes from the device + // buffer into the existing mesh's vertices-heap region (count unchanged) + // and update the root leaf bounds, preserving the mesh handle so + // RTInstance::accelerationStructureReference stays valid. The device + // counterpart of Mesh::RefitProcedural. + __attribute__((import_module("env"), import_name("wgpuRefitMeshBLASDeviceAabbs"))) + extern "C" void wgpuRefitMeshBLASDeviceAabbs( + std::uint32_t meshHandle, + WebGPUBufferRef aabbBufferHandle, std::int32_t count, + float minX, float minY, float minZ, + float maxX, float maxY, float maxZ); + // RT pipeline build. The library composes WGSL by concatenating the // traversal library, generated hit-group switches, and the user- // supplied raygen / miss / closesthit / anyhit bodies. `bindings` is -- 2.54.0 From 8fefb2caeff19b2532fbbc023745abf1ca6cefd4 Mon Sep 17 00:00:00 2001 From: catbot Date: Tue, 16 Jun 2026 14:12:12 +0000 Subject: [PATCH 2/3] test(rt): headless device-buffer BuildProcedural + refit (#37) Extend BLASBuildOptions with a device-local AABB buffer filled via a staging copy + barrier (standing in for a GPU compute producer) fed into the new device-address BuildProcedural/RefitProcedural. Asserts the build/refit succeed, the in-place UPDATE keeps the AS handle + blasAddr, the host aabbBuffer is never allocated (proof of zero-copy), and the validation layer reports no errors. Co-Authored-By: Claude Opus 4.8 --- tests/BLASBuildOptions/main.cpp | 80 +++++++++++++++++++++++++++++++++ 1 file changed, 80 insertions(+) diff --git a/tests/BLASBuildOptions/main.cpp b/tests/BLASBuildOptions/main.cpp index f8e842e..f4fed5f 100644 --- a/tests/BLASBuildOptions/main.cpp +++ b/tests/BLASBuildOptions/main.cpp @@ -197,6 +197,86 @@ int main() { Check(proc.accelerationStructure == procHandleBefore && proc.blasAddr == procAddrBefore, "RefitProcedural kept the same AS handle + blasAddr (in-place UPDATE)"); + // ── Procedural (AABB) BLAS from a DEVICE buffer (issue #37). The boxes + // live in a device-local buffer a GPU pass would write — here filled + // via a staging copy + barrier, standing in for a compute dispatch — + // and fed straight into BuildProcedural/RefitProcedural by device + // address, with no host memcpy through Mesh::aabbBuffer. ──────────── + constexpr std::uint32_t kDevCount = 3; + // Device-local AABB build input: the producer's buffer. Needs the AS + // build-input + device-address usage bits, plus TRANSFER_DST so the + // staging upload (our stand-in GPU writer) can fill it. + VulkanBuffer devAabb; + devAabb.Resize( + VK_BUFFER_USAGE_2_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR + | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT + | VK_BUFFER_USAGE_TRANSFER_DST_BIT, + VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, kDevCount); + VulkanBuffer devStaging; + // SHADER_DEVICE_ADDRESS because VulkanBuffer::Create always queries the + // buffer device address (the staging buffer's address itself is unused). + devStaging.Resize(VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT, kDevCount); + + auto uploadBoxes = [&](VkCommandBuffer cmd, std::span boxes) { + std::memcpy(devStaging.value, boxes.data(), boxes.size() * sizeof(RTAabb)); + devStaging.FlushDevice(); + VkBufferCopy region { .srcOffset = 0, .dstOffset = 0, .size = devAabb.size }; + vkCmdCopyBuffer(cmd, devStaging.buffer, devAabb.buffer, 1, ®ion); + // Order the producing write before the AS build reads the buffer — + // exactly the barrier a real compute producer would need. + VkBufferMemoryBarrier barrier { + .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER, + .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT, + .dstAccessMask = VK_ACCESS_ACCELERATION_STRUCTURE_READ_BIT_KHR, + .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .buffer = devAabb.buffer, .offset = 0, .size = VK_WHOLE_SIZE, + }; + vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_TRANSFER_BIT, + VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR, + 0, 0, nullptr, 1, &barrier, 0, nullptr); + }; + + Mesh devProc; + { + std::array boxes {{ + { .min = {-1,-1,-1}, .max = {1,1,1} }, + { .min = { 2, 2, 2}, .max = {3,3,3} }, + { .min = {-3,-3,-3}, .max = {-2,-2,-2} }, + }}; + VkCommandBuffer cmd = BeginCmd(); + uploadBoxes(cmd, boxes); + devProc.BuildProcedural(devAabb.address, kDevCount, /*opaque*/ false, cmd, + RTBuildOptions{ .preference = RTBuildPreference::FastTrace, .allowUpdate = true }); + SubmitWait(cmd); + } + Check(devProc.blasAddr != 0, "device-buffer BuildProcedural produced a non-zero blasAddr"); + Check(devProc.builtPrimitiveCount == kDevCount, "device-buffer BLAS reports 3 primitives"); + // The host-side aabbBuffer must be untouched — proof there was no memcpy. + Check(devProc.aabbBuffer.buffer == VK_NULL_HANDLE, + "device-buffer path never allocated the host aabbBuffer (zero-copy)"); + + const VkDeviceAddress devAddrBefore = devProc.blasAddr; + const VkAccelerationStructureKHR devHandleBefore = devProc.accelerationStructure; + { + // Same count, moved boxes (re-written into the same device buffer) → + // in-place UPDATE straight from the device address. + std::array boxes {{ + { .min = {-2,-2,-2}, .max = {2,2,2} }, + { .min = { 4, 4, 4}, .max = {5,5,5} }, + { .min = {-5,-5,-5}, .max = {-3,-3,-3} }, + }}; + VkCommandBuffer cmd = BeginCmd(); + uploadBoxes(cmd, boxes); + devProc.RefitProcedural(devAabb.address, kDevCount, cmd); + SubmitWait(cmd); + } + Check(devProc.accelerationStructure == devHandleBefore && devProc.blasAddr == devAddrBefore, + "device-buffer RefitProcedural kept the same AS handle + blasAddr (in-place UPDATE)"); + Check(devProc.aabbBuffer.buffer == VK_NULL_HANDLE, + "device-buffer refit still never touched the host aabbBuffer"); + Check(Device::validationErrorCount == 0, std::format("no Vulkan validation errors ({} seen)", Device::validationErrorCount)); -- 2.54.0 From 0d1312ac0933e2c91df1ab42001c73af56a463a0 Mon Sep 17 00:00:00 2001 From: catbot Date: Tue, 16 Jun 2026 14:12:12 +0000 Subject: [PATCH 3/3] example(RTVolume): GPU-written device-buffer procedural AABBs (#37) A compute shader (aabbs.comp.{glsl,wgsl}) writes the procedural box into a device buffer that BuildProcedural consumes by device address/handle with no host copy; the WebGPU path also refits from it each frame so the box breathes. Verified rendering on both Vulkan (native) and WebGPU (browser). Co-Authored-By: Claude Opus 4.8 --- examples/RTVolume/aabbs.comp.glsl | 40 +++++++++++ examples/RTVolume/aabbs.comp.wgsl | 29 ++++++++ examples/RTVolume/main.cpp | 111 ++++++++++++++++++++++++++---- examples/RTVolume/project.cpp | 6 ++ 4 files changed, 173 insertions(+), 13 deletions(-) create mode 100644 examples/RTVolume/aabbs.comp.glsl create mode 100644 examples/RTVolume/aabbs.comp.wgsl diff --git a/examples/RTVolume/aabbs.comp.glsl b/examples/RTVolume/aabbs.comp.glsl new file mode 100644 index 0000000..0ad513d --- /dev/null +++ b/examples/RTVolume/aabbs.comp.glsl @@ -0,0 +1,40 @@ +#version 460 +#extension GL_EXT_buffer_reference : enable +#extension GL_EXT_scalar_block_layout : enable + +// Issue #37: GPU-resident procedural AABBs. This compute shader is the +// "producer" — it writes the per-box bounding boxes straight into a device +// buffer that Mesh::BuildProcedural / RefitProcedural then consume by device +// address, with no host copy. Here it emits a single unit box [-1,1]^3 (the +// bounding volume the intersection shader turns into a sphere); `time` lets +// the box breathe so a per-frame refit has something to track. A real +// particle system would write one box per particle from simulation state. +// +// The buffer is reached via a buffer_reference (its device address pushed in +// `pc`), so this dispatch needs no descriptor-heap binding for the output. + +layout(buffer_reference, scalar) buffer AabbRef { + float data[]; // 6 floats per box: min.xyz, max.xyz (RTAabb layout) +}; + +layout(push_constant) uniform PC { + AabbRef boxes; + uint count; + float time; +} pc; + +layout(local_size_x = 64) in; + +void main() { + uint i = gl_GlobalInvocationID.x; + if (i >= pc.count) return; + + float r = 1.0 + 0.15 * sin(pc.time); // gently pulsing radius + uint b = i * 6u; + pc.boxes.data[b + 0u] = -r; + pc.boxes.data[b + 1u] = -r; + pc.boxes.data[b + 2u] = -r; + pc.boxes.data[b + 3u] = r; + pc.boxes.data[b + 4u] = r; + pc.boxes.data[b + 5u] = r; +} diff --git a/examples/RTVolume/aabbs.comp.wgsl b/examples/RTVolume/aabbs.comp.wgsl new file mode 100644 index 0000000..aa64f44 --- /dev/null +++ b/examples/RTVolume/aabbs.comp.wgsl @@ -0,0 +1,29 @@ +// Issue #37: GPU producer (WebGPU) for the procedural AABB build input. +// Writes the per-box bounding boxes straight into a device storage buffer +// that Mesh::BuildProcedural / RefitProcedural then consume by handle, with +// no host copy. Emits a single unit box [-r,r]^3 (the bounding volume the +// intersection shader turns into a sphere); `time` lets it breathe so the +// per-frame refit has something to track. A real particle system would write +// one box per particle from simulation state. +// +// PlainComputeShader layout (no rayQuery → user groups start at 1): +// @group(0) @binding(0) uniform Params +// @group(1) @binding(0) storage read_write boxes : array (6 f32/box) + +struct Params { count: u32, time: f32, _0: u32, _1: u32 }; +@group(0) @binding(0) var pc : Params; +@group(1) @binding(0) var boxes : array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + if (i >= pc.count) { return; } + let r = 1.0 + 0.15 * sin(pc.time); + let b = i * 6u; + boxes[b + 0u] = -r; + boxes[b + 1u] = -r; + boxes[b + 2u] = -r; + boxes[b + 3u] = r; + boxes[b + 4u] = r; + boxes[b + 5u] = r; +} diff --git a/examples/RTVolume/main.cpp b/examples/RTVolume/main.cpp index 082494e..8ea2c13 100644 --- a/examples/RTVolume/main.cpp +++ b/examples/RTVolume/main.cpp @@ -88,14 +88,62 @@ int main() { PipelineRTVulkan pipeline; pipeline.Init(cmd, raygenGroups, missGroups, hitGroups, shaderTable); - // ── One procedural unit-box BLAS. The intersection shader treats the - // box as the bounding volume of a radius-1 sphere centred at the - // object origin. opaque=false so the any-hit cut-out runs. ───────── - std::array boxes {{ - { .min = {-1.0f, -1.0f, -1.0f}, .max = {1.0f, 1.0f, 1.0f} }, - }}; + // ── One procedural unit-box BLAS, built from a DEVICE buffer the GPU + // writes (issue #37). A compute shader (aabbs.comp.glsl) writes the + // box [-r,r]^3 into a device-local buffer; that buffer is fed straight + // into Mesh::BuildProcedural by device address — no host copy of the + // AABBs. The intersection shader then treats the box as the bounding + // volume of a sphere centred at the object origin. opaque=false so the + // any-hit cut-out runs. ────────────────────────────────────────────── + constexpr std::uint32_t kBoxCount = 1; + VulkanBuffer aabbDevice; + aabbDevice.Resize( + VK_BUFFER_USAGE_2_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR + | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT + | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, + VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, kBoxCount); + + // The producer reaches its output through a buffer_reference (device + // address pushed below). It binds no heap descriptors, but a + // descriptor-heap compute pipeline still expects a heap bound for its + // push-data path — the render loop binds it every frame, but the init + // command buffer hasn't, so bind it here before dispatching. + { + VkBindHeapInfoEXT resourceHeapInfo { + .sType = VK_STRUCTURE_TYPE_BIND_HEAP_INFO_EXT, + .heapRange = { + .address = descriptorHeap.resourceHeap[window.currentBuffer].address, + .size = static_cast(descriptorHeap.resourceHeap[window.currentBuffer].size), + }, + .reservedRangeOffset = (descriptorHeap.resourceHeap[window.currentBuffer].size - Device::descriptorHeapProperties.minResourceHeapReservedRange) & ~(Device::descriptorHeapProperties.imageDescriptorAlignment - 1), + .reservedRangeSize = Device::descriptorHeapProperties.minResourceHeapReservedRange, + }; + Device::vkCmdBindResourceHeapEXT(cmd, &resourceHeapInfo); + } + + ComputeShader aabbWriter; + aabbWriter.Load("aabbs.comp.spv"); + struct AabbPush { VkDeviceAddress boxes; std::uint32_t count; float time; } push { + aabbDevice.address, kBoxCount, 0.0f }; + aabbWriter.Dispatch(cmd, &push, sizeof(push), (kBoxCount + 63) / 64); + + // Order the producing writes before the BLAS build reads the buffer. + VkMemoryBarrier aabbBarrier { + .sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER, + .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT, + .dstAccessMask = VK_ACCESS_ACCELERATION_STRUCTURE_READ_BIT_KHR, + }; + vkCmdPipelineBarrier(cmd, + VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, + VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR, + 0, 1, &aabbBarrier, 0, nullptr, 0, nullptr); + Mesh sphere; - sphere.BuildProcedural(boxes, /*opaque*/ false, cmd); + // allowUpdate so the GPU-written boxes can be refit in place each frame + // (the per-frame companion to this build — see the headless + // BLASBuildOptions test for the device-buffer refit path). + sphere.BuildProcedural(aabbDevice.address, kBoxCount, /*opaque*/ false, cmd, + RTBuildOptions{ .allowUpdate = true }); // ── Instance grid. ───────────────────────────────────────────────── static std::vector renderers; @@ -241,14 +289,37 @@ int main() { PipelineRTWebGPU pipeline; pipeline.Init(cmd, raygenGroups, missGroups, hitGroups, sbt, bindings); - // ── One procedural unit-box BLAS. The intersection shader treats the - // box as the bounding volume of a radius-1 sphere centred at the - // object origin. opaque=false so the any-hit cut-out runs. ───────── - static std::array boxes {{ - { .min = {-1.0f, -1.0f, -1.0f}, .max = {1.0f, 1.0f, 1.0f} }, + // ── One procedural unit-box BLAS, built from a DEVICE buffer the GPU + // writes (issue #37). A compute shader (aabbs.comp.wgsl) writes the + // box [-r,r]^3 into a device storage buffer; that buffer is fed + // straight into Mesh::BuildProcedural by handle — no host copy of the + // AABBs — and RefitProcedural re-consumes it each frame so the boxes + // can breathe. The intersection shader treats the box as the bounding + // volume of a sphere; opaque=false so the any-hit cut-out runs. ────── + constexpr std::uint32_t kBoxCount = 1; + // worldBounds must enclose every box the producer can write (r ≤ 1.15). + constexpr RTAabb kWorldBounds { .min = {-1.2f, -1.2f, -1.2f}, .max = {1.2f, 1.2f, 1.2f} }; + + static WebGPUBuffer aabbBuf; + aabbBuf.Create(kBoxCount); + + // GPU producer: writes the boxes into aabbBuf. No rayQuery → the storage + // buffer binding lives at group 1. + struct AabbParams { std::uint32_t count; float time; std::uint32_t _0, _1; }; + std::array aabbBindings {{ + { .group = 1, .binding = 0, .kind = UICustomBindingKind::BufferReadWrite, .pushOffset = 0 }, }}; + static PlainComputeShader aabbWriter; + aabbWriter.Load(fs::path("aabbs.comp.wgsl"), sizeof(AabbParams), aabbBindings); + static std::array aabbHandles { aabbBuf.handle }; + + { + AabbParams params { kBoxCount, 0.0f, 0, 0 }; + aabbWriter.Dispatch(¶ms, sizeof(params), aabbHandles, (kBoxCount + 63u) / 64u); + } + static Mesh sphere; - sphere.BuildProcedural(boxes, /*opaque*/ false, cmd); + sphere.BuildProcedural(aabbBuf.handle, kBoxCount, kWorldBounds, /*opaque*/ false, cmd); // ── Camera buffer + handle array. ───────────────────────────────── WebGPUBuffer cameraBuf; @@ -352,6 +423,20 @@ int main() { cameraBuf.FlushDevice(); }); + // ── Per-frame procedural refit from the device buffer (issue #37). The + // GPU producer rewrites the breathing box into aabbBuf each frame and + // RefitProcedural re-consumes it by handle — no host copy, and the + // BLAS handle (blasAddr) stays stable so the TLAS instances built + // above keep referencing it. worldBounds is constant so the TLAS need + // not rebuild. ─────────────────────────────────────────────────────── + EventListener aabbTick(&window.onBeforeUpdate, [&]() { + static float t = 0.0f; + t += kDt; + AabbParams params { kBoxCount, t, 0, 0 }; + aabbWriter.Dispatch(¶ms, sizeof(params), aabbHandles, (kBoxCount + 63u) / 64u); + sphere.RefitProcedural(aabbBuf.handle, kBoxCount, kWorldBounds); + }); + window.Render(); window.StartUpdate(); window.StartSync(); diff --git a/examples/RTVolume/project.cpp b/examples/RTVolume/project.cpp index 0dc138f..154ce76 100644 --- a/examples/RTVolume/project.cpp +++ b/examples/RTVolume/project.cpp @@ -42,6 +42,9 @@ extern "C" Configuration CrafterBuildProject(std::span a cfg.files.emplace_back(fs::path("closesthit.wgsl")); cfg.files.emplace_back(fs::path("miss.wgsl")); cfg.files.emplace_back(fs::path("resolve.wgsl")); + // GPU producer for the AABB build input (issue #37): writes the + // procedural boxes into a device buffer fed straight to BuildProcedural. + cfg.files.emplace_back(fs::path("aabbs.comp.wgsl")); EnableWasiBrowserRuntime(cfg); } else { cfg.shaders.emplace_back(fs::path("raygen.glsl"), std::string("main"), ShaderType::RayGen); @@ -49,6 +52,9 @@ extern "C" Configuration CrafterBuildProject(std::span a cfg.shaders.emplace_back(fs::path("closesthit.glsl"), std::string("main"), ShaderType::ClosestHit); cfg.shaders.emplace_back(fs::path("anyhit.glsl"), std::string("main"), ShaderType::AnyHit); cfg.shaders.emplace_back(fs::path("intersection.glsl"), std::string("main"), ShaderType::Intersect); + // GPU producer for the AABB build input (issue #37): writes the + // procedural boxes into a device buffer fed straight to BuildProcedural. + cfg.shaders.emplace_back(fs::path("aabbs.comp.glsl"), std::string("main"), ShaderType::Compute); } return cfg; } -- 2.54.0