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