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 <noreply@anthropic.com>

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
-    //    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<RTAabb, 1> 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<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
-    //    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<RTAabb, 1> 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<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;
 }