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) => {

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;
+}

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<f32>  (6 f32/box)
+
+struct Params { count: u32, time: f32, _0: u32, _1: u32 };
+@group(0) @binding(0) var<uniform> pc : Params;
+@group(1) @binding(0) var<storage, read_write> boxes : array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    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;
+}

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
+    // ── One procedural unit-box BLAS, built from a DEVICE buffer the GPU
-    //    box as the bounding volume of a radius-1 sphere centred at the
+    //    writes (issue #37). A compute shader (aabbs.comp.glsl) writes the
-    //    object origin. opaque=false so the any-hit cut-out runs. ─────────
+    //    box [-r,r]^3 into a device-local buffer; that buffer is fed straight
-    std::array<RTAabb, 1> boxes {{
+    //    into Mesh::BuildProcedural by device address — no host copy of the
-        { .min = {-1.0f, -1.0f, -1.0f}, .max = {1.0f, 1.0f, 1.0f} },
+    //    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<RTAabb, false> 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<std::uint32_t>(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<RenderingElement3D> 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
+    // ── One procedural unit-box BLAS, built from a DEVICE buffer the GPU
-    //    box as the bounding volume of a radius-1 sphere centred at the
+    //    writes (issue #37). A compute shader (aabbs.comp.wgsl) writes the
-    //    object origin. opaque=false so the any-hit cut-out runs. ─────────
+    //    box [-r,r]^3 into a device storage buffer; that buffer is fed
-    static std::array<RTAabb, 1> boxes {{
+    //    straight into Mesh::BuildProcedural by handle — no host copy of the
-        { .min = {-1.0f, -1.0f, -1.0f}, .max = {1.0f, 1.0f, 1.0f} },
+    //    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<RTAabb, false> 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<UICustomBinding, 1> 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<std::uint32_t, 1> aabbHandles { aabbBuf.handle };
+
+    {
+        AabbParams params { kBoxCount, 0.0f, 0, 0 };
+        aabbWriter.Dispatch(&params, 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<CameraGPU, true> 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<void> aabbTick(&window.onBeforeUpdate, [&]() {
+        static float t = 0.0f;
+        t += kDt;
+        AabbParams params { kBoxCount, t, 0, 0 };
+        aabbWriter.Dispatch(&params, sizeof(params), aabbHandles, (kBoxCount + 63u) / 64u);
+        sphere.RefitProcedural(aabbBuf.handle, kBoxCount, kWorldBounds);
+    });
+
     window.Render();
     window.StartUpdate();
     window.StartSync();

examples/RTVolume/project.cpp

@@ -42,6 +42,9 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> 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<const std::string_view> 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;
 }

implementations/Crafter.Graphics-Mesh-WebGPU.cpp

@@ -342,6 +342,28 @@ void Mesh::BuildProcedural(std::span<const RTAabb> aabbs,
         /*primCount*/ static_cast<std::int32_t>(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<std::int32_t>(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<Vector<float, 3, 3>> vertices,
                  std::span<std::uint32_t>       indices,
                  WebGPUCommandEncoderRef        cmd) {
@@ -359,3 +381,19 @@ void Mesh::RefitProcedural(std::span<const RTAabb> 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<std::uint32_t>(blasAddr), aabbBuffer,
+        static_cast<std::int32_t>(count),
+        worldBounds.min[0], worldBounds.min[1], worldBounds.min[2],
+        worldBounds.max[0], worldBounds.max[1], worldBounds.max[2]);
+}

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
+    // Record an AABB BLAS build (fresh or in-place refit) from a device
-    // in-place refit). Shared by BuildProcedural and RefitProcedural; the
+    // address that already holds the VkAabbPositionsKHR-compatible boxes.
-    // geometry's opaque bit is read from self.opaque so an UPDATE keeps the
+    // Geometry-source agnostic: the host-upload path (RecordProceduralBuild)
-    // exact geometry description of the original build.
+    // points this at self.aabbBuffer, while the device-buffer overloads
-    void RecordProceduralBuild(Mesh& self, std::span<const RTAabb> aabbs, VkBuildAccelerationStructureFlagsKHR flags, bool update, VkCommandBuffer cmd) {
+    // point it straight at the producer's buffer — no copy involved either
-        // 24-byte-stride VkAabbPositionsKHR-compatible build input
+    // way. The geometry's opaque bit is read from self.opaque so an UPDATE
-        // (static_assert'd in the interface). Same usage set as the triangle
+    // keeps the exact geometry description of the original build.
-        // inputs: AS-build read-only + device address. A refit reuses the
+    void RecordProceduralBuildFromAddress(Mesh& self, VkDeviceAddress aabbAddress, std::uint32_t count, std::uint32_t stride, VkBuildAccelerationStructureFlagsKHR flags, bool update, VkCommandBuffer cmd) {
-        // 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<std::uint32_t>(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);
-
         VkAccelerationStructureGeometryAabbsDataKHR aabbsData {
             .sType  = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_AABBS_DATA_KHR,
-            .data   = { .deviceAddress = self.aabbBuffer.address },
+            .data   = { .deviceAddress = aabbAddress },
-            .stride = sizeof(RTAabb)
+            .stride = stride
         };
         VkAccelerationStructureGeometryDataKHR geometryData;
         geometryData.aabbs = aabbsData;
@@ -299,7 +291,25 @@ namespace {
             .flags        = self.opaque ? static_cast<VkGeometryFlagsKHR>(VK_GEOMETRY_OPAQUE_BIT_KHR) : VkGeometryFlagsKHR{}
         };
 
-        RecordBLASBuildFromGeometry(self, blasGeometry, static_cast<std::uint32_t>(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<const RTAabb> 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<std::uint32_t>(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<std::uint32_t>(aabbs.size()), sizeof(RTAabb), flags, update, cmd);
     }
 }
 
@@ -311,6 +321,20 @@ void Mesh::BuildProcedural(std::span<const RTAabb> 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<Vector<float, 3, 3>> verticies, std::span<std::uint32_t> 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<const RTAabb> 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);
+}
+

interfaces/Crafter.Graphics-Mesh.cppm

@@ -125,6 +125,23 @@ export namespace Crafter {
         // exactly like a triangle BLAS.
         void BuildProcedural(std::span<const RTAabb> 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<Vector<float, 3, 3>> verticies, std::span<std::uint32_t> indicies, VkCommandBuffer cmd);
         // Procedural analog of Refit: new object-space boxes, same count.
         void RefitProcedural(std::span<const RTAabb> 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<const RTAabb> 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

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

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<RTAabb, false> 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<RTAabb, true> 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<const RTAabb> 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, &region);
+        // 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<RTAabb, kDevCount> 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<RTAabb, kDevCount> 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));