Support TLAS binding arrays

Cargo.toml

@@ -21,7 +21,7 @@ glam = { version = "0.30", features = ["mint"] }
 gltf = { version = "1.1", default-features = false }
 log = "0.4"
 mint = "0.5"
-naga = { version = "27.0", features = ["wgsl-in", "termcolor"] }
+naga = { git = "https://github.com/gfx-rs/wgpu", rev = "760032a20d6d742d1df28136f36a673ab2584929", features = ["wgsl-in", "termcolor"] }
 profiling = "1"
 slab = "0.4"
 strum = { version = "0.27", features = ["derive"] }

blade-graphics/src/derive.rs

@@ -28,6 +28,9 @@ impl<'a, const N: ResourceIndex> HasShaderBinding for &'a super::TextureArray<N>
 impl HasShaderBinding for super::AccelerationStructure {
     const TYPE: ShaderBinding = ShaderBinding::AccelerationStructure;
 }
+impl<'a, const N: ResourceIndex> HasShaderBinding for &'a super::AccelerationStructureArray<N> {
+    const TYPE: ShaderBinding = ShaderBinding::AccelerationStructureArray { count: N };
+}
 
 pub trait HasVertexAttribute {
     const FORMAT: VertexFormat;

blade-graphics/src/gles/command.rs

@@ -79,6 +79,13 @@
         }
     }
 }
+impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable
+    for &'a crate::AccelerationStructureArray<N>
+{
+    fn bind_to(&self, _ctx: &mut super::PipelineContext, _index: u32) {
+        unimplemented!()
+    }
+}
 
 impl super::CommandEncoder {
     fn begin_pass(&mut self, label: &str) {
@@ -310,7 +317,7 @@
             .push(super::Command::SetViewport(viewport.clone()));
     }
 
     fn set_stencil_reference(&mut self, reference: u32) {
         unimplemented!()
     }
 }
@@ -474,7 +481,7 @@
             .push(super::Command::SetViewport(viewport.clone()));
     }
 
     fn set_stencil_reference(&mut self, reference: u32) {
         unimplemented!()
     }
 }
@@ -833,10 +840,10 @@
                 gl.bind_buffer(glow::PIXEL_UNPACK_BUFFER, None);
             }
             Self::CopyTextureToBuffer {
                 ref src,
                 ref dst,
                 bytes_per_row,
                 ref size,
             } => unimplemented!(),
             Self::ResetFramebuffer => {
                 for &attachment in COLOR_ATTACHMENTS.iter() {
@@ -1060,9 +1067,9 @@
                 gl.scissor(rect.x, rect.y, rect.w as i32, rect.h as i32);
             }
             Self::SetStencilFunc {
                 face,
                 function,
                 reference,
                 read_mask,
             } => unimplemented!(),
             Self::SetStencilOps {

blade-graphics/src/gles/mod.rs

@@ -444,6 +444,7 @@ struct ExecutionContext {
 impl Context {
     pub fn capabilities(&self) -> crate::Capabilities {
         crate::Capabilities {
+            binding_array: false,
             ray_query: crate::ShaderVisibility::empty(),
             sample_count_mask: 0x1 | 0x4, //TODO: accurate info
             dual_source_blending: false,

blade-graphics/src/gles/pipeline.rs

@@ -76,7 +76,10 @@ impl super::Context {
                         }
                         crate::ShaderBinding::TextureArray { .. }
                         | crate::ShaderBinding::BufferArray { .. }
-                        | crate::ShaderBinding::AccelerationStructure => unimplemented!(),
+                        | crate::ShaderBinding::AccelerationStructure
+                        | crate::ShaderBinding::AccelerationStructureArray { .. } => {
+                            unimplemented!()
+                        }
                         crate::ShaderBinding::Plain { .. } => {
                             num_buffers += 1;
                             num_buffers - 1
@@ -248,7 +251,8 @@ impl super::Context {
                         }
                         crate::ShaderBinding::TextureArray { .. }
                         | crate::ShaderBinding::BufferArray { .. }
-                        | crate::ShaderBinding::AccelerationStructure => {
+                        | crate::ShaderBinding::AccelerationStructure
+                        | crate::ShaderBinding::AccelerationStructureArray { .. } => {
                             unimplemented!()
                         }
                         crate::ShaderBinding::Plain { size } => {

blade-graphics/src/lib.rs

@@ -122,7 +122,9 @@ impl From<PlatformError> for NotSupportedError {
 
 #[derive(Clone, Debug, Default, PartialEq)]
 pub struct Capabilities {
-    /// Which shader stages support ray queries
+    /// Support binding arrays of handles.
+    pub binding_array: bool,
+    /// Which shader stages support ray queries.
     pub ray_query: ShaderVisibility,
     /// Bit mask of supported MSAA sample counts.
     pub sample_count_mask: u32,
@@ -281,6 +283,8 @@ impl<T, const N: ResourceIndex> std::ops::IndexMut<ResourceIndex> for ResourceAr
 }
 pub type BufferArray<const N: ResourceIndex> = ResourceArray<BufferPiece, N>;
 pub type TextureArray<const N: ResourceIndex> = ResourceArray<TextureView, N>;
+pub type AccelerationStructureArray<const N: ResourceIndex> =
+    ResourceArray<AccelerationStructure, N>;
 
 #[derive(Clone, Copy, Debug)]
 pub struct TexturePiece {
@@ -644,6 +648,7 @@ pub enum ShaderBinding {
     Buffer,
     BufferArray { count: u32 },
     AccelerationStructure,
+    AccelerationStructureArray { count: u32 },
     Plain { size: u32 },
 }
 

blade-graphics/src/metal/command.rs

@@ -102,6 +102,13 @@
         }
     }
 }
+impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable
+    for &'a crate::AccelerationStructureArray<N>
+{
+    fn bind_to(&self, _ctx: &mut super::PipelineContext, _index: u32) {
+        unimplemented!()
+    }
+}
 
 impl super::TimingData {
     fn add(&mut self, label: &str) -> usize {
@@ -224,7 +231,7 @@
         targets: crate::RenderTargetSet,
     ) -> super::RenderCommandEncoder<'_> {
         let raw = objc2::rc::autoreleasepool(|_| {
             let descriptor = unsafe { metal::MTLRenderPassDescriptor::new() };
 
             for (i, rt) in targets.colors.iter().enumerate() {
                 let at_descriptor =
@@ -362,7 +369,7 @@
         }
 
         let queue = self.queue.lock().unwrap();
         self.raw = Some(objc2::rc::autoreleasepool(|_| unsafe {
             use metal::MTLCommandQueue as _;
             let cmd_buf = queue.commandBufferWithUnretainedReferences().unwrap();
             if !self.name.is_empty() {

blade-graphics/src/metal/mod.rs

@@ -519,6 +519,7 @@
         let device = self.device.lock().unwrap();
 
         crate::Capabilities {
+            binding_array: false,
             ray_query: if device.supportsFamily(metal::MTLGPUFamily::Apple6) {
                 crate::ShaderVisibility::all()
             } else if device.supportsFamily(metal::MTLGPUFamily::Mac2)
@@ -566,7 +567,7 @@
             };
             for i in 0..desc.buffer_count {
                 let label = format!("{}/counter{}", desc.name, i);
                 let sample_buffer = unsafe {
                     csb_desc.setLabel(&objc2_foundation::NSString::from_str(&label));
                     self.device
                         .lock()

blade-graphics/src/metal/pipeline.rs

@@ -76,7 +76,7 @@
     read_mask: u32,
     write_mask: u32,
 ) -> Retained<metal::MTLStencilDescriptor> {
     let desc = unsafe { metal::MTLStencilDescriptor::new() };
     desc.setStencilCompareFunction(super::map_compare_function(face.compare));
     desc.setReadMask(read_mask);
     desc.setWriteMask(write_mask);
@@ -89,7 +89,7 @@
 fn create_depth_stencil_desc(
     state: &crate::DepthStencilState,
 ) -> Retained<metal::MTLDepthStencilDescriptor> {
     let desc = unsafe { metal::MTLDepthStencilDescriptor::new() };
     desc.setDepthCompareFunction(super::map_compare_function(state.depth_compare));
     desc.setDepthWriteEnabled(state.depth_write_enabled);
 
@@ -157,7 +157,8 @@
                     num_buffers - 1
                 }
                 crate::ShaderBinding::TextureArray { .. }
-                | crate::ShaderBinding::BufferArray { .. } => unimplemented!(),
+                | crate::ShaderBinding::BufferArray { .. }
+                | crate::ShaderBinding::AccelerationStructureArray { .. } => unimplemented!(),
                 crate::ShaderBinding::AccelerationStructure => {
                     num_buffers += 1;
                     num_buffers - 1
@@ -268,7 +269,8 @@
                         ..Default::default()
                     },
                     crate::ShaderBinding::TextureArray { .. }
-                    | crate::ShaderBinding::BufferArray { .. } => todo!(),
+                    | crate::ShaderBinding::BufferArray { .. }
+                    | crate::ShaderBinding::AccelerationStructureArray { .. } => todo!(),
                 };
                 naga_resources.resources.insert(res_binding, bind_target);
             }
@@ -446,7 +448,7 @@
                 None
             };
 
             let vertex_descriptor = unsafe { metal::MTLVertexDescriptor::new() };
             for (i, vf) in desc.vertex_fetches.iter().enumerate() {
                 unsafe {
                     let buffer_desc = vertex_descriptor.layouts().objectAtIndexedSubscript(i);

blade-graphics/src/shader.rs

@@ -28,7 +28,15 @@ impl super::Context {
         let flags = naga::valid::ValidationFlags::all() ^ naga::valid::ValidationFlags::BINDINGS;
         let mut caps = naga::valid::Capabilities::empty();
         caps.set(
-            naga::valid::Capabilities::RAY_QUERY | naga::valid::Capabilities::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
+            naga::valid::Capabilities::STORAGE_BUFFER_BINDING_ARRAY
+                | naga::valid::Capabilities::TEXTURE_AND_SAMPLER_BINDING_ARRAY
+                | naga::valid::Capabilities::TEXTURE_AND_SAMPLER_BINDING_ARRAY_NON_UNIFORM_INDEXING
+                | naga::valid::Capabilities::STORAGE_BUFFER_BINDING_ARRAY_NON_UNIFORM_INDEXING,
+            device_caps.binding_array,
+        );
+        caps.set(
+            naga::valid::Capabilities::RAY_QUERY
+                | naga::valid::Capabilities::ACCELERATION_STRUCTURE_BINDING_ARRAY,
             !device_caps.ray_query.is_empty(),
         );
         caps.set(
@@ -170,6 +178,7 @@ impl super::Shader {
                             let count = match proto_binding {
                                 crate::ShaderBinding::TextureArray { count } => count,
                                 crate::ShaderBinding::BufferArray { count } => count,
+                                crate::ShaderBinding::AccelerationStructureArray { count } => count,
                                 _ => 0,
                             };
                             let proto = match module.types[base].inner {
@@ -179,6 +188,9 @@ impl super::Shader {
                                 naga::TypeInner::Struct { .. } => {
                                     crate::ShaderBinding::BufferArray { count }
                                 }
+                                naga::TypeInner::AccelerationStructure { .. } => {
+                                    crate::ShaderBinding::AccelerationStructureArray { count }
+                                }
                                 ref other => panic!("Unsupported binding array for {:?}", other),
                             };
                             (proto, var_access)
@@ -275,6 +287,7 @@ impl super::Shader {
                             interpolation: None,
                             sampling: None,
                             blend_src: None,
+                            per_primitive: false,
                         };
                         for (buffer_index, vertex_fetch) in fetch_states.iter().enumerate() {
                             for (attribute_index, &(at_name, _)) in
@@ -307,6 +320,7 @@ impl super::Shader {
                                 interpolation: None,
                                 sampling: None,
                                 blend_src: None,
+                                per_primitive: false,
                             });
                             location += 1;
                         }

blade-graphics/src/vulkan/command.rs

@@ -130,6 +130,21 @@ impl crate::ShaderBindable for super::AccelerationStructure {
         ctx.write(index, self.raw);
     }
 }
+impl<'a, const N: crate::ResourceIndex> crate::ShaderBindable
+    for &'a crate::AccelerationStructureArray<N>
+{
+    fn bind_to(&self, ctx: &mut super::PipelineContext, index: u32) {
+        assert!(self.data.len() <= N as usize);
+        ctx.write_array(
+            index,
+            self.data
+                .iter()
+                .map(|accel| accel.raw)
+                .cycle()
+                .take(N as usize),
+        );
+    }
+}
 
 impl crate::TexturePiece {
     fn subresource_layers(&self) -> vk::ImageSubresourceLayers {

blade-graphics/src/vulkan/init.rs

@@ -44,6 +44,7 @@ struct AdapterCapabilities {
     device_information: crate::DeviceInformation,
     queue_family_index: u32,
     layered: bool,
+    binding_array: bool,
     ray_tracing: Option<RayTracingCapabilities>,
     buffer_marker: bool,
     shader_info: bool,
@@ -209,6 +210,14 @@ unsafe fn inspect_adapter(
         true
     };
 
+    let supports_descriptor_indexing = api_version >= vk::API_VERSION_1_2
+        || supported_extensions.contains(&vk::EXT_DESCRIPTOR_INDEXING_NAME);
+    let binding_array = supports_descriptor_indexing
+        && descriptor_indexing_features.descriptor_binding_partially_bound == vk::TRUE
+        && descriptor_indexing_features.shader_storage_buffer_array_non_uniform_indexing
+            == vk::TRUE
+        && descriptor_indexing_features.shader_sampled_image_array_non_uniform_indexing == vk::TRUE;
+
     let ray_tracing = if !supported_extensions.contains(&vk::KHR_ACCELERATION_STRUCTURE_NAME)
         || !supported_extensions.contains(&vk::KHR_RAY_QUERY_NAME)
     {
@@ -277,6 +286,7 @@ unsafe fn inspect_adapter(
         device_information,
         queue_family_index,
         layered: portability_subset_properties.min_vertex_input_binding_stride_alignment != 0,
+        binding_array,
         ray_tracing,
         buffer_marker,
         shader_info,
@@ -457,9 +467,13 @@ impl super::Context {
                 log::info!("Enabling Vulkan Portability");
                 device_extensions.push(vk::KHR_PORTABILITY_SUBSET_NAME);
             }
+            let needs_descriptor_indexing =
+                capabilities.binding_array || capabilities.ray_tracing.is_some();
+            if needs_descriptor_indexing && capabilities.api_version < vk::API_VERSION_1_2 {
+                device_extensions.push(vk::EXT_DESCRIPTOR_INDEXING_NAME);
+            }
             if capabilities.ray_tracing.is_some() {
                 if capabilities.api_version < vk::API_VERSION_1_2 {
-                    device_extensions.push(vk::EXT_DESCRIPTOR_INDEXING_NAME);
                     device_extensions.push(vk::KHR_BUFFER_DEVICE_ADDRESS_NAME);
                     device_extensions.push(vk::KHR_SHADER_FLOAT_CONTROLS_NAME);
                     device_extensions.push(vk::KHR_SPIRV_1_4_NAME);
@@ -514,13 +528,16 @@ impl super::Context {
             let mut khr_buffer_device_address;
             let mut khr_acceleration_structure;
             let mut khr_ray_query;
-            if capabilities.ray_tracing.is_some() {
+            if needs_descriptor_indexing {
                 ext_descriptor_indexing = vk::PhysicalDeviceDescriptorIndexingFeaturesEXT {
                     shader_storage_buffer_array_non_uniform_indexing: vk::TRUE,
                     shader_sampled_image_array_non_uniform_indexing: vk::TRUE,
                     descriptor_binding_partially_bound: vk::TRUE,
                     ..Default::default()
                 };
+                device_create_info = device_create_info.push_next(&mut ext_descriptor_indexing);
+            }
+            if capabilities.ray_tracing.is_some() {
                 khr_buffer_device_address = vk::PhysicalDeviceBufferDeviceAddressFeaturesKHR {
                     buffer_device_address: vk::TRUE,
                     ..Default::default()
@@ -534,7 +551,6 @@ impl super::Context {
                     ..Default::default()
                 };
                 device_create_info = device_create_info
-                    .push_next(&mut ext_descriptor_indexing)
                     .push_next(&mut khr_buffer_device_address)
                     .push_next(&mut khr_acceleration_structure)
                     .push_next(&mut khr_ray_query);
@@ -747,6 +763,7 @@ impl super::Context {
                     .limits
                     .framebuffer_depth_sample_counts,
             dual_source_blending: capabilities.dual_source_blending,
+            binding_array: capabilities.binding_array,
             instance,
             entry,
         })
@@ -766,6 +783,7 @@ impl super::Context {
 
     pub fn capabilities(&self) -> crate::Capabilities {
         crate::Capabilities {
+            binding_array: self.binding_array,
             ray_query: match self.device.ray_tracing {
                 Some(_) => crate::ShaderVisibility::all(),
                 None => crate::ShaderVisibility::empty(),

blade-graphics/src/vulkan/mod.rs

@@ -159,6 +159,7 @@ pub struct Context {
     min_buffer_alignment: u64,
     sample_count_flags: vk::SampleCountFlags,
     dual_source_blending: bool,
+    binding_array: bool,
     instance: Instance,
     entry: ash::Entry,
 }

blade-graphics/src/vulkan/pipeline.rs

@@ -20,7 +20,8 @@ impl super::Context {
             for (binding_index, &(_, binding)) in layout.bindings.iter().enumerate() {
                 let binding_array_size = match binding {
                     crate::ShaderBinding::TextureArray { count }
-                    | crate::ShaderBinding::BufferArray { count } => Some(count),
+                    | crate::ShaderBinding::BufferArray { count }
+                    | crate::ShaderBinding::AccelerationStructureArray { count } => Some(count),
                     _ => None,
                 };
                 let rb = naga::ResourceBinding {
@@ -51,6 +52,7 @@ impl super::Context {
             bounds_check_policies: naga::proc::BoundsCheckPolicies::default(),
             zero_initialize_workgroup_memory: spv::ZeroInitializeWorkgroupMemoryMode::None,
             force_loop_bounding: false,
+            ray_query_initialization_tracking: true,
             use_storage_input_output_16: false,
             debug_info: None,
         }
@@ -241,6 +243,12 @@ impl super::Context {
                     1u32,
                     vk::DescriptorBindingFlags::empty(),
                 ),
+                crate::ShaderBinding::AccelerationStructureArray { count } => (
+                    vk::DescriptorType::ACCELERATION_STRUCTURE_KHR,
+                    mem::size_of::<vk::AccelerationStructureKHR>(),
+                    count,
+                    vk::DescriptorBindingFlags::PARTIALLY_BOUND,
+                ),
                 crate::ShaderBinding::Plain { size } => (
                     vk::DescriptorType::INLINE_UNIFORM_BLOCK_EXT,
                     1,

blade-render/code/fill-gbuf.wgsl

@@ -1,3 +1,4 @@
+enable wgpu_ray_query;
 #include "quaternion.inc.wgsl"
 #include "camera.inc.wgsl"
 #include "debug.inc.wgsl"

blade-render/code/ray-trace.wgsl

@@ -1,3 +1,4 @@
+enable wgpu_ray_query;
 #include "quaternion.inc.wgsl"
 #include "random.inc.wgsl"
 #include "env-importance.inc.wgsl"

docs/CHANGELOG.md

@@ -1,5 +1,9 @@
 Changelog for *Blade* project
 
+## blade-graphics-0.8 (TBD)
+- graphics
+  - separate `Capabilities` flag for binding arrays, including TLAS arrays
+
 ## blade-graphics-0.7 (27 Sep 2025)
 
 - graphics

examples/ray-query/shader.wgsl

@@ -1,3 +1,4 @@
+enable wgpu_ray_query;
 const MAX_BOUNCES: i32 = 3;
 
 struct Parameters {