Support getting hit vertex positions (#7183)

CHANGELOG.md

@@ -148,6 +148,8 @@ By @jamienicol in [#6929](https://github.com/gfx-rs/wgpu/pull/6929) and [#7080](
 - Added `CommandEncoder::transition_resources()` for native API interop, and allowing users to slightly optimize barriers. By @JMS55 in [#6678](https://github.com/gfx-rs/wgpu/pull/6678).
 - Add `wgpu_hal::vulkan::Adapter::texture_format_as_raw` for native API interop. By @JMS55 in [#7228](https://github.com/gfx-rs/wgpu/pull/7228).
 
+- Support getting vertices of the hit triangle when raytracing. By @Vecvec in [#7183](https://github.com/gfx-rs/wgpu/pull/7183) .
+
 #### Naga
 
 - Add support for unsigned types when calling textureLoad with the level parameter. By @ygdrasil-io in [#7058](https://github.com/gfx-rs/wgpu/pull/7058).

docs/api-specs/ray_tracing.md

@@ -50,6 +50,8 @@ Naming is mostly taken from vulkan.
 ```wgsl
 // - Initializes the `ray_query` to check where (if anywhere) the ray defined by `ray_desc` hits in `acceleration_structure
 rayQueryInitialize(rq: ptr<function, ray_query>, acceleration_structure: acceleration_structure, ray_desc: RayDesc)
+// Overload.
+rayQueryInitialize(rq: ptr<function, ray_query<vertex_return>>, acceleration_structure: acceleration_structure<vertex_return>, ray_desc: RayDesc)
 
 // - Traces the ray in the initialized ray_query (partially) through the scene.
 // - Returns true if a triangle that was hit by the ray was in a `Blas` that is not marked as opaque.
@@ -63,6 +65,8 @@ rayQueryInitialize(rq: ptr<function, ray_query>, acceleration_structure: acceler
 // - Calling this function multiple times will cause the ray traversal to continue if it was interrupted by a `Candidate`
 //   intersection.
 rayQueryProceed(rq: ptr<function, ray_query>) -> bool
+// Overload.
+rayQueryProceed(rq: ptr<function, ray_query<vertex_return>>) -> bool
 
 // - Generates a hit from procedural geometry at a particular distance.
 rayQueryGenerateIntersection(hit_t: f32)
@@ -75,9 +79,19 @@ rayQueryTerminate()
 
 // - Returns intersection details about a hit considered `Committed`.
 rayQueryGetCommittedIntersection(rq: ptr<function, ray_query>) -> RayIntersection
+// Overload.
+rayQueryGetCommittedIntersection(rq: ptr<function, ray_query<vertex_return>>) -> RayIntersection
 
 // - Returns intersection details about a hit considered `Candidate`.
 rayQueryGetCandidateIntersection(rq: ptr<function, ray_query>) -> RayIntersection
+// Overload.
+rayQueryGetCandidateIntersection(rq: ptr<function, ray_query<vertex_return>>) -> RayIntersection
+
+// - Returns the vertices of the hit triangle considered `Committed`.
+getCommittedHitVertexPositions(rq: ptr<function, ray_query<vertex_return>>) -> array<vec3<f32>, 3>
+
+// - Returns the vertices of the hit triangle considered `Candidate`.
+getCandidateHitVertexPositions(rq: ptr<function, ray_query<vertex_return>>) -> array<vec3<f32>, 3>
 ```
 
 > [!CAUTION]
@@ -89,6 +103,11 @@ rayQueryGetCandidateIntersection(rq: ptr<function, ray_query>) -> RayIntersectio
 >   `Candidate`.
 > - Calling `rayQueryGetCandidateIntersection` when `rayQueryProceed`'s latest return on this ray query is considered
 >   `Committed`.
+> - Calling `getCommittedHitVertexPositions` when `rayQueryProceed`'s latest return on this ray query is considered
+>   `Candidate`.
+> - Calling `getCandidateHitVertexPositions` when `rayQueryProceed`'s latest return on this ray query is considered
+>   `Committed`.
+> - Calling `get*HitVertexPositions` when the last `rayQueryProceed` did not hit a triangle
 > - Calling `rayQueryProceed` when `rayQueryInitialize` has not previously been called on this ray query
 > - Calling `rayQueryGenerateIntersection` on a query with last intersection kind not being
 >   `RAY_QUERY_INTERSECTION_AABB`,

examples/features/src/lib.rs

@@ -17,6 +17,7 @@ pub mod msaa_line;
 pub mod multiple_render_targets;
 pub mod ray_cube_compute;
 pub mod ray_cube_fragment;
+pub mod ray_cube_normals;
 pub mod ray_scene;
 pub mod ray_shadows;
 pub mod ray_traced_triangle;

examples/features/src/main.rs

@@ -170,6 +170,12 @@ const EXAMPLES: &[ExampleDesc] = &[
         webgl: false,
         webgpu: false,
     },
+    ExampleDesc {
+        name: "ray_cube_normals",
+        function: wgpu_examples::ray_cube_normals::main,
+        webgl: false,  // No Ray-tracing extensions
+        webgpu: false, // No Ray-tracing extensions (yet)
+    },
 ];
 
 fn get_example_name() -> Option<String> {

examples/features/src/ray_cube_normals/README.md

@@ -0,0 +1,14 @@
+# ray-cube
+
+This example renders a ray traced cube with hardware acceleration.
+A separate compute shader is used to perform the ray queries. 
+
+## To Run
+
+```
+cargo run --bin wgpu-examples ray_cube_normals
+```
+
+## Screenshots
+
+![Cube example](screenshot.png)

examples/features/src/ray_cube_normals/blit.wgsl

@@ -0,0 +1,52 @@
+struct VertexOutput {
+    @builtin(position) position: vec4<f32>,
+    @location(0) tex_coords: vec2<f32>,
+};
+
+// meant to be called with 3 vertex indices: 0, 1, 2
+// draws one large triangle over the clip space like this:
+// (the asterisks represent the clip space bounds)
+//-1,1           1,1
+// ---------------------------------
+// |              *              .
+// |              *           .
+// |              *        .
+// |              *      .
+// |              *    . 
+// |              * .
+// |***************
+// |            . 1,-1 
+// |          .
+// |       .
+// |     .
+// |   .
+// |.
+@vertex
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
+    var result: VertexOutput;
+    let x = i32(vertex_index) / 2;
+    let y = i32(vertex_index) & 1;
+    let tc = vec2<f32>(
+        f32(x) * 2.0,
+        f32(y) * 2.0
+    );
+    result.position = vec4<f32>(
+        tc.x * 2.0 - 1.0,
+        1.0 - tc.y * 2.0,
+        0.0, 1.0
+    );
+    result.tex_coords = tc;
+    return result;
+}
+
+@group(0)
+@binding(0)
+var r_color: texture_2d<f32>;
+@group(0)
+@binding(1)
+var r_sampler: sampler;
+
+@fragment
+fn fs_main(vertex: VertexOutput) -> @location(0) vec4<f32> {
+    return textureSample(r_color, r_sampler, vertex.tex_coords);
+}

examples/features/src/ray_cube_normals/mod.rs

@@ -0,0 +1,492 @@
+use std::{borrow::Cow, future::Future, iter, mem, pin::Pin, task, time::Instant};
+
+use bytemuck::{Pod, Zeroable};
+use glam::{Affine3A, Mat4, Quat, Vec3};
+use wgpu::util::DeviceExt;
+
+use wgpu::StoreOp;
+
+// from cube
+#[repr(C)]
+#[derive(Clone, Copy, Pod, Zeroable)]
+struct Vertex {
+    _pos: [f32; 4],
+    _tex_coord: [f32; 2],
+}
+
+fn vertex(pos: [i8; 3], tc: [i8; 2]) -> Vertex {
+    Vertex {
+        _pos: [pos[0] as f32, pos[1] as f32, pos[2] as f32, 1.0],
+        _tex_coord: [tc[0] as f32, tc[1] as f32],
+    }
+}
+
+fn create_vertices() -> (Vec<Vertex>, Vec<u16>) {
+    let vertex_data = [
+        // top (0, 0, 1)
+        vertex([-1, -1, 1], [0, 0]),
+        vertex([1, -1, 1], [1, 0]),
+        vertex([1, 1, 1], [1, 1]),
+        vertex([-1, 1, 1], [0, 1]),
+        // bottom (0, 0, -1)
+        vertex([-1, 1, -1], [1, 0]),
+        vertex([1, 1, -1], [0, 0]),
+        vertex([1, -1, -1], [0, 1]),
+        vertex([-1, -1, -1], [1, 1]),
+        // right (1, 0, 0)
+        vertex([1, -1, -1], [0, 0]),
+        vertex([1, 1, -1], [1, 0]),
+        vertex([1, 1, 1], [1, 1]),
+        vertex([1, -1, 1], [0, 1]),
+        // left (-1, 0, 0)
+        vertex([-1, -1, 1], [1, 0]),
+        vertex([-1, 1, 1], [0, 0]),
+        vertex([-1, 1, -1], [0, 1]),
+        vertex([-1, -1, -1], [1, 1]),
+        // front (0, 1, 0)
+        vertex([1, 1, -1], [1, 0]),
+        vertex([-1, 1, -1], [0, 0]),
+        vertex([-1, 1, 1], [0, 1]),
+        vertex([1, 1, 1], [1, 1]),
+        // back (0, -1, 0)
+        vertex([1, -1, 1], [0, 0]),
+        vertex([-1, -1, 1], [1, 0]),
+        vertex([-1, -1, -1], [1, 1]),
+        vertex([1, -1, -1], [0, 1]),
+    ];
+
+    let index_data: &[u16] = &[
+        0, 1, 2, 2, 3, 0, // top
+        4, 5, 6, 6, 7, 4, // bottom
+        8, 9, 10, 10, 11, 8, // right
+        12, 13, 14, 14, 15, 12, // left
+        16, 17, 18, 18, 19, 16, // front
+        20, 21, 22, 22, 23, 20, // back
+    ];
+
+    (vertex_data.to_vec(), index_data.to_vec())
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Pod, Zeroable)]
+struct Uniforms {
+    view_inverse: Mat4,
+    proj_inverse: Mat4,
+}
+
+#[inline]
+fn affine_to_rows(mat: &Affine3A) -> [f32; 12] {
+    let row_0 = mat.matrix3.row(0);
+    let row_1 = mat.matrix3.row(1);
+    let row_2 = mat.matrix3.row(2);
+    let translation = mat.translation;
+    [
+        row_0.x,
+        row_0.y,
+        row_0.z,
+        translation.x,
+        row_1.x,
+        row_1.y,
+        row_1.z,
+        translation.y,
+        row_2.x,
+        row_2.y,
+        row_2.z,
+        translation.z,
+    ]
+}
+
+/// A wrapper for `pop_error_scope` futures that panics if an error occurs.
+///
+/// Given a future `inner` of an `Option<E>` for some error type `E`,
+/// wait for the future to be ready, and panic if its value is `Some`.
+///
+/// This can be done simpler with `FutureExt`, but we don't want to add
+/// a dependency just for this small case.
+struct ErrorFuture<F> {
+    inner: F,
+}
+impl<F: Future<Output = Option<wgpu::Error>>> Future for ErrorFuture<F> {
+    type Output = ();
+    fn poll(self: Pin<&mut Self>, cx: &mut task::Context<'_>) -> task::Poll<()> {
+        let inner = unsafe { self.map_unchecked_mut(|me| &mut me.inner) };
+        inner.poll(cx).map(|error| {
+            if let Some(e) = error {
+                panic!("Rendering {}", e);
+            }
+        })
+    }
+}
+
+struct Example {
+    rt_target: wgpu::Texture,
+    tlas_package: wgpu::TlasPackage,
+    compute_pipeline: wgpu::ComputePipeline,
+    compute_bind_group: wgpu::BindGroup,
+    blit_pipeline: wgpu::RenderPipeline,
+    blit_bind_group: wgpu::BindGroup,
+    start_inst: Instant,
+}
+
+impl crate::framework::Example for Example {
+    // Don't want srgb, so normals show up better.
+    const SRGB: bool = false;
+    fn required_features() -> wgpu::Features {
+        wgpu::Features::EXPERIMENTAL_RAY_QUERY
+            | wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
+            | wgpu::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN
+    }
+
+    fn required_downlevel_capabilities() -> wgpu::DownlevelCapabilities {
+        wgpu::DownlevelCapabilities::default()
+    }
+    fn required_limits() -> wgpu::Limits {
+        wgpu::Limits::default()
+    }
+
+    fn init(
+        config: &wgpu::SurfaceConfiguration,
+        _adapter: &wgpu::Adapter,
+        device: &wgpu::Device,
+        queue: &wgpu::Queue,
+    ) -> Self {
+        let side_count = 8;
+
+        let rt_target = device.create_texture(&wgpu::TextureDescriptor {
+            label: Some("rt_target"),
+            size: wgpu::Extent3d {
+                width: config.width,
+                height: config.height,
+                depth_or_array_layers: 1,
+            },
+            mip_level_count: 1,
+            sample_count: 1,
+            dimension: wgpu::TextureDimension::D2,
+            format: wgpu::TextureFormat::Rgba8Unorm,
+            usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::STORAGE_BINDING,
+            view_formats: &[wgpu::TextureFormat::Rgba8Unorm],
+        });
+
+        let rt_view = rt_target.create_view(&wgpu::TextureViewDescriptor {
+            label: None,
+            format: Some(wgpu::TextureFormat::Rgba8Unorm),
+            dimension: Some(wgpu::TextureViewDimension::D2),
+            usage: None,
+            aspect: wgpu::TextureAspect::All,
+            base_mip_level: 0,
+            mip_level_count: None,
+            base_array_layer: 0,
+            array_layer_count: None,
+        });
+
+        let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
+            label: Some("rt_sampler"),
+            address_mode_u: wgpu::AddressMode::ClampToEdge,
+            address_mode_v: wgpu::AddressMode::ClampToEdge,
+            address_mode_w: wgpu::AddressMode::ClampToEdge,
+            mag_filter: wgpu::FilterMode::Linear,
+            min_filter: wgpu::FilterMode::Linear,
+            mipmap_filter: wgpu::FilterMode::Nearest,
+            ..Default::default()
+        });
+
+        let uniforms = {
+            let view = Mat4::look_at_rh(Vec3::new(0.0, 0.0, 2.5), Vec3::ZERO, Vec3::Y);
+            let proj = Mat4::perspective_rh(
+                59.0_f32.to_radians(),
+                config.width as f32 / config.height as f32,
+                0.001,
+                1000.0,
+            );
+
+            Uniforms {
+                view_inverse: view.inverse(),
+                proj_inverse: proj.inverse(),
+            }
+        };
+
+        let uniform_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("Uniform Buffer"),
+            contents: bytemuck::cast_slice(&[uniforms]),
+            usage: wgpu::BufferUsages::UNIFORM,
+        });
+
+        let (vertex_data, index_data) = create_vertices();
+
+        let vertex_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("Vertex Buffer"),
+            contents: bytemuck::cast_slice(&vertex_data),
+            usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::BLAS_INPUT,
+        });
+
+        let index_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("Index Buffer"),
+            contents: bytemuck::cast_slice(&index_data),
+            usage: wgpu::BufferUsages::INDEX | wgpu::BufferUsages::BLAS_INPUT,
+        });
+
+        let blas_geo_size_desc = wgpu::BlasTriangleGeometrySizeDescriptor {
+            vertex_format: wgpu::VertexFormat::Float32x3,
+            vertex_count: vertex_data.len() as u32,
+            index_format: Some(wgpu::IndexFormat::Uint16),
+            index_count: Some(index_data.len() as u32),
+            flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE,
+        };
+
+        let blas = device.create_blas(
+            &wgpu::CreateBlasDescriptor {
+                label: None,
+                flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE
+                    | wgpu::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN,
+                update_mode: wgpu::AccelerationStructureUpdateMode::Build,
+            },
+            wgpu::BlasGeometrySizeDescriptors::Triangles {
+                descriptors: vec![blas_geo_size_desc.clone()],
+            },
+        );
+
+        let tlas = device.create_tlas(&wgpu::CreateTlasDescriptor {
+            label: None,
+            flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE
+                | wgpu::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN,
+            update_mode: wgpu::AccelerationStructureUpdateMode::Build,
+            max_instances: side_count * side_count,
+        });
+
+        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("rt_computer"),
+            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("shader.wgsl"))),
+        });
+
+        let blit_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("blit"),
+            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("blit.wgsl"))),
+        });
+
+        let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("rt"),
+            layout: None,
+            module: &shader,
+            entry_point: None,
+            compilation_options: Default::default(),
+            cache: None,
+        });
+
+        let compute_bind_group_layout = compute_pipeline.get_bind_group_layout(0);
+
+        let compute_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &compute_bind_group_layout,
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::TextureView(&rt_view),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: uniform_buf.as_entire_binding(),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 2,
+                    resource: wgpu::BindingResource::AccelerationStructure(&tlas),
+                },
+            ],
+        });
+
+        let blit_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+            label: Some("blit"),
+            layout: None,
+            vertex: wgpu::VertexState {
+                module: &blit_shader,
+                entry_point: Some("vs_main"),
+                compilation_options: Default::default(),
+                buffers: &[],
+            },
+            fragment: Some(wgpu::FragmentState {
+                module: &blit_shader,
+                entry_point: Some("fs_main"),
+                compilation_options: Default::default(),
+                targets: &[Some(config.format.into())],
+            }),
+            primitive: wgpu::PrimitiveState {
+                topology: wgpu::PrimitiveTopology::TriangleList,
+                ..Default::default()
+            },
+            depth_stencil: None,
+            multisample: wgpu::MultisampleState::default(),
+            multiview: None,
+            cache: None,
+        });
+
+        let blit_bind_group_layout = blit_pipeline.get_bind_group_layout(0);
+
+        let blit_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &blit_bind_group_layout,
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::TextureView(&rt_view),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: wgpu::BindingResource::Sampler(&sampler),
+                },
+            ],
+        });
+
+        let mut tlas_package = wgpu::TlasPackage::new(tlas);
+
+        let dist = 3.0;
+
+        for x in 0..side_count {
+            for y in 0..side_count {
+                tlas_package[(x + y * side_count) as usize] = Some(wgpu::TlasInstance::new(
+                    &blas,
+                    affine_to_rows(&Affine3A::from_rotation_translation(
+                        Quat::from_rotation_y(45.9_f32.to_radians()),
+                        Vec3 {
+                            x: x as f32 * dist,
+                            y: y as f32 * dist,
+                            z: -30.0,
+                        },
+                    )),
+                    0,
+                    0xff,
+                ));
+            }
+        }
+
+        let mut encoder =
+            device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+
+        encoder.build_acceleration_structures(
+            iter::once(&wgpu::BlasBuildEntry {
+                blas: &blas,
+                geometry: wgpu::BlasGeometries::TriangleGeometries(vec![
+                    wgpu::BlasTriangleGeometry {
+                        size: &blas_geo_size_desc,
+                        vertex_buffer: &vertex_buf,
+                        first_vertex: 0,
+                        vertex_stride: mem::size_of::<Vertex>() as u64,
+                        index_buffer: Some(&index_buf),
+                        first_index: Some(0),
+                        transform_buffer: None,
+                        transform_buffer_offset: None,
+                    },
+                ]),
+            }),
+            iter::once(&tlas_package),
+        );
+
+        queue.submit(Some(encoder.finish()));
+
+        let start_inst = Instant::now();
+
+        Example {
+            rt_target,
+            tlas_package,
+            compute_pipeline,
+            compute_bind_group,
+            blit_pipeline,
+            blit_bind_group,
+            start_inst,
+        }
+    }
+
+    fn update(&mut self, _event: winit::event::WindowEvent) {
+        //empty
+    }
+
+    fn resize(
+        &mut self,
+        _config: &wgpu::SurfaceConfiguration,
+        _device: &wgpu::Device,
+        _queue: &wgpu::Queue,
+    ) {
+    }
+
+    fn render(&mut self, view: &wgpu::TextureView, device: &wgpu::Device, queue: &wgpu::Queue) {
+        device.push_error_scope(wgpu::ErrorFilter::Validation);
+
+        let anim_time = self.start_inst.elapsed().as_secs_f64() as f32;
+
+        self.tlas_package[0].as_mut().unwrap().transform =
+            affine_to_rows(&Affine3A::from_rotation_translation(
+                Quat::from_euler(
+                    glam::EulerRot::XYZ,
+                    anim_time * 0.342,
+                    anim_time * 0.254,
+                    anim_time * 0.832,
+                ),
+                Vec3 {
+                    x: 0.0,
+                    y: 0.0,
+                    z: -6.0,
+                },
+            ));
+
+        let mut encoder =
+            device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+
+        encoder.build_acceleration_structures(iter::empty(), iter::once(&self.tlas_package));
+
+        {
+            let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: None,
+                timestamp_writes: None,
+            });
+            cpass.set_pipeline(&self.compute_pipeline);
+            cpass.set_bind_group(0, Some(&self.compute_bind_group), &[]);
+            cpass.dispatch_workgroups(self.rt_target.width() / 8, self.rt_target.height() / 8, 1);
+        }
+
+        {
+            let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+                label: None,
+                color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+                    view,
+                    resolve_target: None,
+                    ops: wgpu::Operations {
+                        load: wgpu::LoadOp::Clear(wgpu::Color::GREEN),
+                        store: StoreOp::Store,
+                    },
+                })],
+                depth_stencil_attachment: None,
+                timestamp_writes: None,
+                occlusion_query_set: None,
+            });
+
+            rpass.set_pipeline(&self.blit_pipeline);
+            rpass.set_bind_group(0, Some(&self.blit_bind_group), &[]);
+            rpass.draw(0..3, 0..1);
+        }
+
+        queue.submit(Some(encoder.finish()));
+    }
+}
+
+pub fn main() {
+    crate::framework::run::<Example>("ray-cube");
+}
+
+#[cfg(test)]
+#[wgpu_test::gpu_test]
+static TEST: crate::framework::ExampleTestParams = crate::framework::ExampleTestParams {
+    name: "ray_cube_normals",
+    image_path: "/examples/src/ray_cube_normals/screenshot.png",
+    width: 1024,
+    height: 768,
+    optional_features: wgpu::Features::default(),
+    base_test_parameters: wgpu_test::TestParameters {
+        required_features: <Example as crate::framework::Example>::required_features(),
+        required_limits: <Example as crate::framework::Example>::required_limits(),
+        force_fxc: false,
+        skips: vec![],
+        failures: Vec::new(),
+        required_downlevel_caps:
+            <Example as crate::framework::Example>::required_downlevel_capabilities(),
+    },
+    comparisons: &[wgpu_test::ComparisonType::Mean(0.02)],
+    _phantom: std::marker::PhantomData::<Example>,
+};

examples/features/src/ray_cube_normals/shader.wgsl

@@ -0,0 +1,85 @@
+/*
+let RAY_FLAG_NONE = 0x00u;
+let RAY_FLAG_OPAQUE = 0x01u;
+let RAY_FLAG_NO_OPAQUE = 0x02u;
+let RAY_FLAG_TERMINATE_ON_FIRST_HIT = 0x04u;
+let RAY_FLAG_SKIP_CLOSEST_HIT_SHADER = 0x08u;
+let RAY_FLAG_CULL_BACK_FACING = 0x10u;
+let RAY_FLAG_CULL_FRONT_FACING = 0x20u;
+let RAY_FLAG_CULL_OPAQUE = 0x40u;
+let RAY_FLAG_CULL_NO_OPAQUE = 0x80u;
+let RAY_FLAG_SKIP_TRIANGLES = 0x100u;
+let RAY_FLAG_SKIP_AABBS = 0x200u;
+
+let RAY_QUERY_INTERSECTION_NONE = 0u;
+let RAY_QUERY_INTERSECTION_TRIANGLE = 1u;
+let RAY_QUERY_INTERSECTION_GENERATED = 2u;
+let RAY_QUERY_INTERSECTION_AABB = 4u;
+
+struct RayDesc {
+    flags: u32,
+    cull_mask: u32,
+    t_min: f32,
+    t_max: f32,
+    origin: vec3<f32>,
+    dir: vec3<f32>,
+}
+
+struct RayIntersection {
+    kind: u32,
+    t: f32,
+    instance_custom_index: u32,
+    instance_id: u32,
+    sbt_record_offset: u32,
+    geometry_index: u32,
+    primitive_index: u32,
+    barycentrics: vec2<f32>,
+    front_face: bool,
+    object_to_world: mat4x3<f32>,
+    world_to_object: mat4x3<f32>,
+}
+*/
+
+struct Uniforms {
+    view_inv: mat4x4<f32>,
+    proj_inv: mat4x4<f32>,
+};
+
+@group(0) @binding(0)
+var output: texture_storage_2d<rgba8unorm, write>;
+
+@group(0) @binding(1)
+var<uniform> uniforms: Uniforms;
+
+@group(0) @binding(2)
+var acc_struct: acceleration_structure<vertex_return>;
+
+@compute @workgroup_size(8, 8)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let target_size = textureDimensions(output);
+    var color =  vec4<f32>(vec2<f32>(global_id.xy) / vec2<f32>(target_size), 0.0, 1.0);
+
+
+	let pixel_center = vec2<f32>(global_id.xy) + vec2<f32>(0.5);
+	let in_uv = pixel_center/vec2<f32>(target_size.xy);
+	let d = in_uv * 2.0 - 1.0;
+
+	let origin = (uniforms.view_inv * vec4<f32>(0.0,0.0,0.0,1.0)).xyz;
+	let temp = uniforms.proj_inv * vec4<f32>(d.x, d.y, 1.0, 1.0);
+	let direction = (uniforms.view_inv * vec4<f32>(normalize(temp.xyz), 0.0)).xyz;
+
+    var rq: ray_query<vertex_return>;
+    rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.1, 200.0, origin, direction));
+    rayQueryProceed(&rq);
+
+    let intersection = rayQueryGetCommittedIntersection(&rq);
+    if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) {
+        var positions : array<vec3f, 3> = getCommittedHitVertexPositions(&rq);
+        // The cube should change colour as it rotates because it's normals are changing
+        let normals = intersection.object_to_world * vec4f(normalize(cross(positions[0] - positions[1], positions[0] - positions[2])), 0.0);
+        // the y is negated because the texture coordinates are inverted
+        color = vec4f(normals.x, -normals.y, normals.z, 1.0);
+    }
+
+    textureStore(output, global_id.xy, color);
+}

examples/features/src/ray_shadows/mod.rs

@@ -192,7 +192,9 @@ impl crate::framework::Example for Example {
                 wgpu::BindGroupLayoutEntry {
                     binding: 1,
                     visibility: wgpu::ShaderStages::FRAGMENT,
-                    ty: wgpu::BindingType::AccelerationStructure,
+                    ty: wgpu::BindingType::AccelerationStructure {
+                        vertex_return: false,
+                    },
                     count: None,
                 },
             ],

examples/features/src/ray_traced_triangle/mod.rs

@@ -72,7 +72,9 @@ impl crate::framework::Example for Example {
                 wgpu::BindGroupLayoutEntry {
                     binding: 2,
                     visibility: wgpu::ShaderStages::COMPUTE,
-                    ty: wgpu::BindingType::AccelerationStructure,
+                    ty: wgpu::BindingType::AccelerationStructure {
+                        vertex_return: false,
+                    },
                     count: None,
                 },
             ],

naga/src/back/dot/mod.rs

@@ -727,6 +727,11 @@ fn write_function_expressions(
             }
             E::SubgroupBallotResult => ("SubgroupBallotResult".into(), 4),
             E::SubgroupOperationResult { .. } => ("SubgroupOperationResult".into(), 4),
+            E::RayQueryVertexPositions { query, committed } => {
+                edges.insert("", query);
+                let ty = if committed { "Committed" } else { "Candidate" };
+                (format!("get{}HitVertexPositions", ty).into(), 4)
+            }
         };
 
         // give uniform expressions an outline

naga/src/back/glsl/mod.rs

@@ -1082,8 +1082,8 @@ impl<'a, W: Write> Writer<'a, W> {
             | TypeInner::Struct { .. }
             | TypeInner::Image { .. }
             | TypeInner::Sampler { .. }
-            | TypeInner::AccelerationStructure
-            | TypeInner::RayQuery
+            | TypeInner::AccelerationStructure { .. }
+            | TypeInner::RayQuery { .. }
             | TypeInner::BindingArray { .. } => {
                 return Err(Error::Custom(format!("Unable to write type {inner:?}")))
             }
@@ -4017,7 +4017,8 @@ impl<'a, W: Write> Writer<'a, W> {
                 write!(self.out, ".length())")?
             }
             // not supported yet
-            Expression::RayQueryGetIntersection { .. } => unreachable!(),
+            Expression::RayQueryGetIntersection { .. }
+            | Expression::RayQueryVertexPositions { .. } => unreachable!(),
         }
 
         Ok(())

naga/src/back/hlsl/ray.rs

@@ -28,7 +28,12 @@ impl<W: Write> super::Writer<'_, W> {
     ) -> BackendResult {
         self.write_type(module, module.special_types.ray_intersection.unwrap())?;
         write!(self.out, " GetCommittedIntersection(")?;
-        self.write_value_type(module, &TypeInner::RayQuery)?;
+        self.write_value_type(
+            module,
+            &TypeInner::RayQuery {
+                vertex_return: false,
+            },
+        )?;
         writeln!(self.out, " rq) {{")?;
         write!(self.out, "    ")?;
         self.write_type(module, module.special_types.ray_intersection.unwrap())?;
@@ -94,7 +99,12 @@ impl<W: Write> super::Writer<'_, W> {
     ) -> BackendResult {
         self.write_type(module, module.special_types.ray_intersection.unwrap())?;
         write!(self.out, " GetCandidateIntersection(")?;
-        self.write_value_type(module, &TypeInner::RayQuery)?;
+        self.write_value_type(
+            module,
+            &TypeInner::RayQuery {
+                vertex_return: false,
+            },
+        )?;
         writeln!(self.out, " rq) {{")?;
         write!(self.out, "    ")?;
         self.write_type(module, module.special_types.ray_intersection.unwrap())?;

naga/src/back/hlsl/writer.rs

@@ -1362,10 +1362,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
             TypeInner::Array { base, size, .. } | TypeInner::BindingArray { base, size } => {
                 self.write_array_size(module, base, size)?;
             }
-            TypeInner::AccelerationStructure => {
+            TypeInner::AccelerationStructure { .. } => {
                 write!(self.out, "RaytracingAccelerationStructure")?;
             }
-            TypeInner::RayQuery => {
+            TypeInner::RayQuery { .. } => {
                 // these are constant flags, there are dynamic flags also but constant flags are not supported by naga
                 write!(self.out, "RayQuery<RAY_FLAG_NONE>")?;
             }
@@ -1540,7 +1540,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
 
             match module.types[local.ty].inner {
                 // from https://microsoft.github.io/DirectX-Specs/d3d/Raytracing.html#tracerayinline-example-1 it seems that ray queries shouldn't be zeroed
-                TypeInner::RayQuery => {}
+                TypeInner::RayQuery { .. } => {}
                 _ => {
                     write!(self.out, " = ")?;
                     // Write the local initializer if needed
@@ -3953,6 +3953,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     write!(self.out, ")")?;
                 }
             }
+            // Not supported yet
+            Expression::RayQueryVertexPositions { .. } => unreachable!(),
             // Nothing to do here, since call expression already cached
             Expression::CallResult(_)
             | Expression::AtomicResult { .. }

naga/src/back/mod.rs

@@ -265,7 +265,7 @@ impl crate::TypeInner {
         match *self {
             crate::TypeInner::Image { .. }
             | crate::TypeInner::Sampler { .. }
-            | crate::TypeInner::AccelerationStructure => true,
+            | crate::TypeInner::AccelerationStructure { .. } => true,
             _ => false,
         }
     }

naga/src/back/msl/writer.rs

@@ -289,10 +289,16 @@ impl Display for TypeContext<'_> {
             crate::TypeInner::Sampler { comparison: _ } => {
                 write!(out, "{NAMESPACE}::sampler")
             }
-            crate::TypeInner::AccelerationStructure => {
+            crate::TypeInner::AccelerationStructure { vertex_return } => {
+                if vertex_return {
+                    unimplemented!("metal does not support vertex ray hit return")
+                }
                 write!(out, "{RT_NAMESPACE}::instance_acceleration_structure")
             }
-            crate::TypeInner::RayQuery => {
+            crate::TypeInner::RayQuery { vertex_return } => {
+                if vertex_return {
+                    unimplemented!("metal does not support vertex ray hit return")
+                }
                 write!(out, "{RAY_QUERY_TYPE}")
             }
             crate::TypeInner::BindingArray { base, .. } => {
@@ -569,8 +575,8 @@ impl crate::Type {
             // handle types may be different, depending on the global var access, so we always inline them
             Ti::Image { .. }
             | Ti::Sampler { .. }
-            | Ti::AccelerationStructure
-            | Ti::RayQuery
+            | Ti::AccelerationStructure { .. }
+            | Ti::RayQuery { .. }
             | Ti::BindingArray { .. } => false,
         }
     }
@@ -2358,6 +2364,9 @@ impl<W: Write> Writer<W> {
                     write!(self.out, ")")?;
                 }
             }
+            crate::Expression::RayQueryVertexPositions { .. } => {
+                unimplemented!()
+            }
             crate::Expression::RayQueryGetIntersection {
                 query,
                 committed: _,
@@ -3815,12 +3824,12 @@ impl<W: Write> Writer<W> {
         let mut uses_ray_query = false;
         for (_, ty) in module.types.iter() {
             match ty.inner {
-                crate::TypeInner::AccelerationStructure => {
+                crate::TypeInner::AccelerationStructure { .. } => {
                     if options.lang_version < (2, 4) {
                         return Err(Error::UnsupportedRayTracing);
                     }
                 }
-                crate::TypeInner::RayQuery => {
+                crate::TypeInner::RayQuery { .. } => {
                     if options.lang_version < (2, 4) {
                         return Err(Error::UnsupportedRayTracing);
                     }

naga/src/back/pipeline_constants.rs

@@ -638,6 +638,12 @@ fn adjust_expr(new_pos: &HandleVec<Expression, Handle<Expression>>, expr: &mut E
         | Expression::WorkGroupUniformLoadResult { ty: _ }
         | Expression::SubgroupBallotResult
         | Expression::SubgroupOperationResult { .. } => {}
+        Expression::RayQueryVertexPositions {
+            ref mut query,
+            committed: _,
+        } => {
+            adjust(query);
+        }
     }
 }
 

naga/src/back/spv/block.rs

@@ -1930,6 +1930,13 @@ impl BlockContext<'_> {
                 ));
                 id
             }
+            crate::Expression::RayQueryVertexPositions { query, committed } => {
+                self.writer.require_any(
+                    "RayQueryVertexPositions",
+                    &[spirv::Capability::RayQueryPositionFetchKHR],
+                )?;
+                self.write_ray_query_return_vertex_position(query, block, committed)
+            }
         };
 
         self.cached[expr_handle] = id;

naga/src/back/spv/instructions.rs

@@ -795,6 +795,20 @@ impl super::Instruction {
         instruction
     }
 
+    pub(super) fn ray_query_return_vertex_position(
+        result_type_id: Word,
+        id: Word,
+        query: Word,
+        intersection: Word,
+    ) -> Self {
+        let mut instruction = Self::new(Op::RayQueryGetIntersectionTriangleVertexPositionsKHR);
+        instruction.set_type(result_type_id);
+        instruction.set_result(id);
+        instruction.add_operand(query);
+        instruction.add_operand(intersection);
+        instruction
+    }
+
     pub(super) fn ray_query_get_intersection(
         op: Op,
         result_type_id: Word,

naga/src/back/spv/mod.rs

@@ -488,8 +488,8 @@ impl LocalType {
                 class,
             } => LocalType::Image(LocalImageType::from_inner(dim, arrayed, class)),
             crate::TypeInner::Sampler { comparison: _ } => LocalType::Sampler,
-            crate::TypeInner::AccelerationStructure => LocalType::AccelerationStructure,
-            crate::TypeInner::RayQuery => LocalType::RayQuery,
+            crate::TypeInner::AccelerationStructure { .. } => LocalType::AccelerationStructure,
+            crate::TypeInner::RayQuery { .. } => LocalType::RayQuery,
             crate::TypeInner::Array { .. }
             | crate::TypeInner::Struct { .. }
             | crate::TypeInner::BindingArray { .. } => return None,

naga/src/back/spv/ray.rs

@@ -116,17 +116,8 @@ impl Writer {
             class: spirv::StorageClass::Function,
         }));
 
-        let rq_ty = ir_module
-            .types
-            .get(&Type {
-                name: None,
-                inner: TypeInner::RayQuery,
-            })
-            .expect("ray_query type should have been populated by the variable passed into this!");
-        let argument_type_id = self.get_type_id(LookupType::Local(LocalType::Pointer {
-            base: rq_ty,
-            class: spirv::StorageClass::Function,
-        }));
+        let argument_type_id = self.get_ray_query_pointer_id(ir_module);
+
         let func_ty = self.get_function_type(LookupFunctionType {
             parameter_type_ids: vec![argument_type_id],
             return_type_id: intersection_type_id,
@@ -626,4 +617,39 @@ impl BlockContext<'_> {
             crate::RayQueryFunction::Terminate => {}
         }
     }
+
+    pub(super) fn write_ray_query_return_vertex_position(
+        &mut self,
+        query: Handle<crate::Expression>,
+        block: &mut Block,
+        is_committed: bool,
+    ) -> spirv::Word {
+        let query_id = self.cached[query];
+        let id = self.gen_id();
+        let result = self
+            .ir_module
+            .special_types
+            .ray_vertex_return
+            .expect("type should have been populated");
+        let intersection_id =
+            self.writer
+                .get_constant_scalar(crate::Literal::U32(if is_committed {
+                    spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR
+                } else {
+                    spirv::RayQueryIntersection::RayQueryCandidateIntersectionKHR
+                } as _));
+        block
+            .body
+            .push(Instruction::ray_query_return_vertex_position(
+                *self
+                    .writer
+                    .lookup_type
+                    .get(&LookupType::Handle(result))
+                    .expect("type should have been populated"),
+                id,
+                query_id,
+                intersection_id,
+            ));
+        id
+    }
 }

naga/src/back/spv/writer.rs

@@ -251,6 +251,30 @@ impl Writer {
         }))
     }
 
+    pub(super) fn get_ray_query_pointer_id(&mut self, module: &crate::Module) -> Word {
+        let rq_ty = module
+            .types
+            .get(&crate::Type {
+                name: None,
+                inner: crate::TypeInner::RayQuery {
+                    vertex_return: false,
+                },
+            })
+            .or_else(|| {
+                module.types.get(&crate::Type {
+                    name: None,
+                    inner: crate::TypeInner::RayQuery {
+                        vertex_return: true,
+                    },
+                })
+            })
+            .expect("ray_query type should have been populated by the variable passed into this!");
+        self.get_type_id(LookupType::Local(LocalType::Pointer {
+            base: rq_ty,
+            class: spirv::StorageClass::Function,
+        }))
+    }
+
     /// Return a SPIR-V type for a pointer to `resolution`.
     ///
     /// The given `resolution` must be one that we can represent
@@ -913,7 +937,7 @@ impl Writer {
                 id,
                 spirv::StorageClass::Function,
                 init_word.or_else(|| match ir_module.types[variable.ty].inner {
-                    crate::TypeInner::RayQuery => None,
+                    crate::TypeInner::RayQuery { .. } => None,
                     _ => {
                         let type_id = context.get_type_id(LookupType::Handle(variable.ty));
                         Some(context.writer.write_constant_null(type_id))
@@ -1138,10 +1162,10 @@ impl Writer {
                     _ => {}
                 }
             }
-            crate::TypeInner::AccelerationStructure => {
+            crate::TypeInner::AccelerationStructure { .. } => {
                 self.require_any("Acceleration Structure", &[spirv::Capability::RayQueryKHR])?;
             }
-            crate::TypeInner::RayQuery => {
+            crate::TypeInner::RayQuery { .. } => {
                 self.require_any("Ray Query", &[spirv::Capability::RayQueryKHR])?;
             }
             crate::TypeInner::Atomic(crate::Scalar { width: 8, kind: _ }) => {
@@ -1324,8 +1348,8 @@ impl Writer {
                 | crate::TypeInner::ValuePointer { .. }
                 | crate::TypeInner::Image { .. }
                 | crate::TypeInner::Sampler { .. }
-                | crate::TypeInner::AccelerationStructure
-                | crate::TypeInner::RayQuery => unreachable!(),
+                | crate::TypeInner::AccelerationStructure { .. }
+                | crate::TypeInner::RayQuery { .. } => unreachable!(),
             };
 
             instruction.to_words(&mut self.logical_layout.declarations);
@@ -2199,9 +2223,13 @@ impl Writer {
             .any(|arg| has_view_index_check(ir_module, arg.binding.as_ref(), arg.ty));
         let mut has_ray_query = ir_module.special_types.ray_desc.is_some()
             | ir_module.special_types.ray_intersection.is_some();
+        let has_vertex_return = ir_module.special_types.ray_vertex_return.is_some();
 
         for (_, &crate::Type { ref inner, .. }) in ir_module.types.iter() {
-            if let &crate::TypeInner::AccelerationStructure | &crate::TypeInner::RayQuery = inner {
+            // spirv does not know whether these have vertex return - that is done by us
+            if let &crate::TypeInner::AccelerationStructure { .. }
+            | &crate::TypeInner::RayQuery { .. } = inner
+            {
                 has_ray_query = true
             }
         }
@@ -2219,6 +2247,10 @@ impl Writer {
             Instruction::extension("SPV_KHR_ray_query")
                 .to_words(&mut self.logical_layout.extensions)
         }
+        if has_vertex_return {
+            Instruction::extension("SPV_KHR_ray_tracing_position_fetch")
+                .to_words(&mut self.logical_layout.extensions);
+        }
         Instruction::type_void(self.void_type).to_words(&mut self.logical_layout.declarations);
         Instruction::ext_inst_import(self.gl450_ext_inst_id, "GLSL.std.450")
             .to_words(&mut self.logical_layout.ext_inst_imports);

naga/src/back/wgsl/writer.rs

@@ -628,7 +628,10 @@ impl<W: Write> Writer<W> {
                 }
                 write!(self.out, ">")?;
             }
-            TypeInner::AccelerationStructure => write!(self.out, "acceleration_structure")?,
+            TypeInner::AccelerationStructure { vertex_return } => {
+                let caps = if vertex_return { "<vertex_return>" } else { "" };
+                write!(self.out, "acceleration_structure{}", caps)?
+            }
             _ => {
                 return Err(Error::Unimplemented(format!("write_value_type {inner:?}")));
             }
@@ -1890,7 +1893,8 @@ impl<W: Write> Writer<W> {
                 write!(self.out, ")")?
             }
             // Not supported yet
-            Expression::RayQueryGetIntersection { .. } => unreachable!(),
+            Expression::RayQueryGetIntersection { .. }
+            | Expression::RayQueryVertexPositions { .. } => unreachable!(),
             // Nothing to do here, since call expression already cached
             Expression::CallResult(_)
             | Expression::AtomicResult { .. }

naga/src/compact/expressions.rs

@@ -185,6 +185,12 @@ impl ExpressionTracer<'_> {
                     Iq::NumLevels | Iq::NumLayers | Iq::NumSamples => {}
                 }
             }
+            Ex::RayQueryVertexPositions {
+                query,
+                committed: _,
+            } => {
+                self.expressions_used.insert(query);
+            }
             Ex::Unary { op: _, expr } => {
                 self.expressions_used.insert(expr);
             }
@@ -402,6 +408,10 @@ impl ModuleMap {
                 ref mut query,
                 committed: _,
             } => adjust(query),
+            Ex::RayQueryVertexPositions {
+                ref mut query,
+                committed: _,
+            } => adjust(query),
         }
     }
 

naga/src/compact/mod.rs

@@ -317,6 +317,7 @@ impl<'module> ModuleTracer<'module> {
         let crate::SpecialTypes {
             ref ray_desc,
             ref ray_intersection,
+            ref ray_vertex_return,
             ref predeclared_types,
         } = *special_types;
 
@@ -326,6 +327,9 @@ impl<'module> ModuleTracer<'module> {
         if let Some(ray_intersection) = *ray_intersection {
             self.types_used.insert(ray_intersection);
         }
+        if let Some(ray_vertex_return) = *ray_vertex_return {
+            self.types_used.insert(ray_vertex_return);
+        }
         for (_, &handle) in predeclared_types {
             self.types_used.insert(handle);
         }
@@ -462,6 +466,7 @@ impl ModuleMap {
         let crate::SpecialTypes {
             ref mut ray_desc,
             ref mut ray_intersection,
+            ref mut ray_vertex_return,
             ref mut predeclared_types,
         } = *special;
 
@@ -472,6 +477,10 @@ impl ModuleMap {
             self.types.adjust(ray_intersection);
         }
 
+        if let Some(ref mut ray_vertex_return) = *ray_vertex_return {
+            self.types.adjust(ray_vertex_return);
+        }
+
         for handle in predeclared_types.values_mut() {
             self.types.adjust(handle);
         }

naga/src/compact/types.rs

@@ -20,8 +20,8 @@ impl TypeTracer<'_> {
             | Ti::ValuePointer { .. }
             | Ti::Image { .. }
             | Ti::Sampler { .. }
-            | Ti::AccelerationStructure
-            | Ti::RayQuery => {}
+            | Ti::AccelerationStructure { .. }
+            | Ti::RayQuery { .. } => {}
 
             // Types that do contain handles.
             Ti::Array {
@@ -75,8 +75,8 @@ impl ModuleMap {
             | Ti::ValuePointer { .. }
             | Ti::Image { .. }
             | Ti::Sampler { .. }
-            | Ti::AccelerationStructure
-            | Ti::RayQuery => {}
+            | Ti::AccelerationStructure { .. }
+            | Ti::RayQuery { .. } => {}
 
             // Types that do contain handles.
             Ti::Pointer {

naga/src/front/type_gen.rs

@@ -104,6 +104,36 @@ impl crate::Module {
         handle
     }
 
+    /// Make sure the types for the vertex return are in the module's type
+    pub fn generate_vertex_return_type(&mut self) -> Handle<crate::Type> {
+        if let Some(handle) = self.special_types.ray_vertex_return {
+            return handle;
+        }
+        let ty_vec3f = self.types.insert(
+            crate::Type {
+                name: None,
+                inner: crate::TypeInner::Vector {
+                    size: crate::VectorSize::Tri,
+                    scalar: crate::Scalar::F32,
+                },
+            },
+            Span::UNDEFINED,
+        );
+        let array = self.types.insert(
+            crate::Type {
+                name: None,
+                inner: crate::TypeInner::Array {
+                    base: ty_vec3f,
+                    size: crate::ArraySize::Constant(core::num::NonZeroU32::new(3).unwrap()),
+                    stride: 16,
+                },
+            },
+            Span::UNDEFINED,
+        );
+        self.special_types.ray_vertex_return = Some(array);
+        array
+    }
+
     /// Populate this module's [`SpecialTypes::ray_intersection`] type.
     ///
     /// [`SpecialTypes::ray_intersection`] is the type of a

naga/src/front/wgsl/lower/conversion.rs

@@ -442,8 +442,8 @@ impl crate::TypeInner {
             | Ti::Struct { .. }
             | Ti::Image { .. }
             | Ti::Sampler { .. }
-            | Ti::AccelerationStructure
-            | Ti::RayQuery
+            | Ti::AccelerationStructure { .. }
+            | Ti::RayQuery { .. }
             | Ti::BindingArray { .. } => None,
         }
     }
@@ -468,8 +468,8 @@ impl crate::TypeInner {
             Ti::Struct { .. }
             | Ti::Image { .. }
             | Ti::Sampler { .. }
-            | Ti::AccelerationStructure
-            | Ti::RayQuery
+            | Ti::AccelerationStructure { .. }
+            | Ti::RayQuery { .. }
             | Ti::BindingArray { .. } => None,
         }
     }

naga/src/front/wgsl/lower/mod.rs

@@ -2724,6 +2724,30 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                                 .push(crate::Statement::RayQuery { query, fun }, span);
                             return Ok(None);
                         }
+                        "getCommittedHitVertexPositions" => {
+                            let mut args = ctx.prepare_args(arguments, 1, span);
+                            let query = self.ray_query_pointer(args.next()?, ctx)?;
+                            args.finish()?;
+
+                            let _ = ctx.module.generate_vertex_return_type();
+
+                            crate::Expression::RayQueryVertexPositions {
+                                query,
+                                committed: true,
+                            }
+                        }
+                        "getCandidateHitVertexPositions" => {
+                            let mut args = ctx.prepare_args(arguments, 1, span);
+                            let query = self.ray_query_pointer(args.next()?, ctx)?;
+                            args.finish()?;
+
+                            let _ = ctx.module.generate_vertex_return_type();
+
+                            crate::Expression::RayQueryVertexPositions {
+                                query,
+                                committed: false,
+                            }
+                        }
                         "rayQueryProceed" => {
                             let mut args = ctx.prepare_args(arguments, 1, span);
                             let query = self.ray_query_pointer(args.next()?, ctx)?;
@@ -3336,8 +3360,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                 class,
             },
             ast::Type::Sampler { comparison } => crate::TypeInner::Sampler { comparison },
-            ast::Type::AccelerationStructure => crate::TypeInner::AccelerationStructure,
-            ast::Type::RayQuery => crate::TypeInner::RayQuery,
+            ast::Type::AccelerationStructure { vertex_return } => {
+                crate::TypeInner::AccelerationStructure { vertex_return }
+            }
+            ast::Type::RayQuery { vertex_return } => crate::TypeInner::RayQuery { vertex_return },
             ast::Type::BindingArray { base, size } => {
                 let base = self.resolve_ast_type(base, ctx)?;
                 let size = self.array_size(size, ctx)?;
@@ -3407,7 +3433,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
 
         match *resolve_inner!(ctx, pointer) {
             crate::TypeInner::Pointer { base, .. } => match ctx.module.types[base].inner {
-                crate::TypeInner::RayQuery => Ok(pointer),
+                crate::TypeInner::RayQuery { .. } => Ok(pointer),
                 ref other => {
                     log::error!("Pointer type to {:?} passed to ray query op", other);
                     Err(Error::InvalidRayQueryPointer(span))

naga/src/front/wgsl/parse/ast.rs

@@ -243,8 +243,12 @@ pub enum Type<'a> {
     Sampler {
         comparison: bool,
     },
-    AccelerationStructure,
-    RayQuery,
+    AccelerationStructure {
+        vertex_return: bool,
+    },
+    RayQuery {
+        vertex_return: bool,
+    },
     RayDesc,
     RayIntersection,
     BindingArray {

naga/src/front/wgsl/parse/lexer.rs

@@ -466,6 +466,28 @@ impl<'a> Lexer<'a> {
         Ok((format, access))
     }
 
+    pub(in crate::front::wgsl) fn next_acceleration_structure_flags(
+        &mut self,
+    ) -> Result<bool, Error<'a>> {
+        Ok(if self.skip(Token::Paren('<')) {
+            if !self.skip(Token::Paren('>')) {
+                let (name, span) = self.next_ident_with_span()?;
+                let ret = if name == "vertex_return" {
+                    true
+                } else {
+                    return Err(Error::UnknownAttribute(span));
+                };
+                self.skip(Token::Separator(','));
+                self.expect(Token::Paren('>'))?;
+                ret
+            } else {
+                false
+            }
+        } else {
+            false
+        })
+    }
+
     pub(in crate::front::wgsl) fn open_arguments(&mut self) -> Result<(), Error<'a>> {
         self.expect(Token::Paren('('))
     }

naga/src/front/wgsl/parse/mod.rs

@@ -1654,8 +1654,14 @@ impl Parser {
                     class: crate::ImageClass::Storage { format, access },
                 }
             }
-            "acceleration_structure" => ast::Type::AccelerationStructure,
-            "ray_query" => ast::Type::RayQuery,
+            "acceleration_structure" => {
+                let vertex_return = lexer.next_acceleration_structure_flags()?;
+                ast::Type::AccelerationStructure { vertex_return }
+            }
+            "ray_query" => {
+                let vertex_return = lexer.next_acceleration_structure_flags()?;
+                ast::Type::RayQuery { vertex_return }
+            }
             "RayDesc" => ast::Type::RayDesc,
             "RayIntersection" => ast::Type::RayIntersection,
             _ => return Ok(None),

naga/src/front/wgsl/to_wgsl.rs

@@ -122,8 +122,14 @@ impl crate::TypeInner {
                 format!("texture{class_suffix}{dim_suffix}{array_suffix}{type_in_brackets}")
             }
             Ti::Sampler { .. } => "sampler".to_string(),
-            Ti::AccelerationStructure => "acceleration_structure".to_string(),
-            Ti::RayQuery => "ray_query".to_string(),
+            Ti::AccelerationStructure { vertex_return } => {
+                let caps = if vertex_return { "<vertex_return>" } else { "" };
+                format!("acceleration_structure{}", caps)
+            }
+            Ti::RayQuery { vertex_return } => {
+                let caps = if vertex_return { "<vertex_return>" } else { "" };
+                format!("ray_query{}", caps)
+            }
             Ti::BindingArray { base, size, .. } => {
                 let member_type = &gctx.types[base];
                 let base = member_type.name.as_deref().unwrap_or("unknown");

naga/src/lib.rs

@@ -847,10 +847,10 @@ pub enum TypeInner {
     Sampler { comparison: bool },
 
     /// Opaque object representing an acceleration structure of geometry.
-    AccelerationStructure,
+    AccelerationStructure { vertex_return: bool },
 
     /// Locally used handle for ray queries.
-    RayQuery,
+    RayQuery { vertex_return: bool },
 
     /// Array of bindings.
     ///
@@ -1697,6 +1697,14 @@ pub enum Expression {
     /// a pointer to a structure containing a runtime array in its' last field.
     ArrayLength(Handle<Expression>),
 
+    /// Get the Positions of the triangle hit by the [`RayQuery`]
+    ///
+    /// [`RayQuery`]: Statement::RayQuery
+    RayQueryVertexPositions {
+        query: Handle<Expression>,
+        committed: bool,
+    },
+
     /// Result of a [`Proceed`] [`RayQuery`] statement.
     ///
     /// [`Proceed`]: RayQueryFunction::Proceed
@@ -2341,6 +2349,11 @@ pub struct SpecialTypes {
     /// this if needed and return the handle.
     pub ray_intersection: Option<Handle<Type>>,
 
+    /// Type for `RayVertexReturn
+    ///
+    /// Call [`Module::generate_vertex_return_type`]
+    pub ray_vertex_return: Option<Handle<Type>>,
+
     /// Types for predeclared wgsl types instantiated on demand.
     ///
     /// Call [`Module::generate_predeclared_type`] to populate this if

naga/src/proc/constant_evaluator.rs

@@ -963,7 +963,9 @@ impl<'a> ConstantEvaluator<'a> {
             Expression::ImageSample { .. }
             | Expression::ImageLoad { .. }
             | Expression::ImageQuery { .. } => Err(ConstantEvaluatorError::ImageExpression),
-            Expression::RayQueryProceedResult | Expression::RayQueryGetIntersection { .. } => {
+            Expression::RayQueryProceedResult
+            | Expression::RayQueryGetIntersection { .. }
+            | Expression::RayQueryVertexPositions { .. } => {
                 Err(ConstantEvaluatorError::RayQueryExpression)
             }
             Expression::SubgroupBallotResult => Err(ConstantEvaluatorError::SubgroupExpression),

naga/src/proc/layouter.rs

@@ -234,8 +234,8 @@ impl Layouter {
                 }
                 Ti::Image { .. }
                 | Ti::Sampler { .. }
-                | Ti::AccelerationStructure
-                | Ti::RayQuery
+                | Ti::AccelerationStructure { .. }
+                | Ti::RayQuery { .. }
                 | Ti::BindingArray { .. } => TypeLayout {
                     size,
                     alignment: Alignment::ONE,

naga/src/proc/type_methods.rs

@@ -164,8 +164,8 @@ impl crate::TypeInner {
             Self::Struct { span, .. } => span,
             Self::Image { .. }
             | Self::Sampler { .. }
-            | Self::AccelerationStructure
-            | Self::RayQuery
+            | Self::AccelerationStructure { .. }
+            | Self::RayQuery { .. }
             | Self::BindingArray { .. } => 0,
         }
     }
@@ -276,8 +276,8 @@ impl crate::TypeInner {
             | crate::TypeInner::Struct { .. }
             | crate::TypeInner::Image { .. }
             | crate::TypeInner::Sampler { .. }
-            | crate::TypeInner::AccelerationStructure
-            | crate::TypeInner::RayQuery
+            | crate::TypeInner::AccelerationStructure { .. }
+            | crate::TypeInner::RayQuery { .. }
             | crate::TypeInner::BindingArray { .. } => None,
         }
     }
@@ -298,8 +298,8 @@ impl crate::TypeInner {
             | crate::TypeInner::Struct { .. }
             | crate::TypeInner::Image { .. }
             | crate::TypeInner::Sampler { .. }
-            | crate::TypeInner::AccelerationStructure
-            | crate::TypeInner::RayQuery
+            | crate::TypeInner::AccelerationStructure { .. }
+            | crate::TypeInner::RayQuery { .. }
             | crate::TypeInner::BindingArray { .. } => false,
         }
     }

naga/src/proc/typifier.rs

@@ -149,6 +149,7 @@ impl Clone for TypeResolution {
                     scalar,
                     space,
                 },
+                Ti::Array { base, size, stride } => Ti::Array { base, size, stride },
                 _ => unreachable!("Unexpected clone type: {:?}", v),
             }),
         }
@@ -900,6 +901,13 @@ impl<'a> ResolveContext<'a> {
                     .ok_or(ResolveError::MissingSpecialType)?;
                 TypeResolution::Handle(result)
             }
+            crate::Expression::RayQueryVertexPositions { .. } => {
+                let result = self
+                    .special_types
+                    .ray_vertex_return
+                    .ok_or(ResolveError::MissingSpecialType)?;
+                TypeResolution::Handle(result)
+            }
             crate::Expression::SubgroupBallotResult => TypeResolution::Value(Ti::Vector {
                 scalar: crate::Scalar::U32,
                 size: crate::VectorSize::Quad,

naga/src/valid/analyzer.rs

@@ -809,6 +809,13 @@ impl FunctionInfo {
                 non_uniform_result: Some(handle),
                 requirements: UniformityRequirements::empty(),
             },
+            E::RayQueryVertexPositions {
+                query,
+                committed: _,
+            } => Uniformity {
+                non_uniform_result: self.add_ref(query),
+                requirements: UniformityRequirements::empty(),
+            },
         };
 
         let ty = resolve_context.resolve(expression, |h| Ok(&self[h].ty))?;

naga/src/valid/expression.rs

@@ -1726,7 +1726,28 @@ impl super::Validator {
                     base,
                     space: crate::AddressSpace::Function,
                 } => match resolver.types[base].inner {
-                    Ti::RayQuery => ShaderStages::all(),
+                    Ti::RayQuery { .. } => ShaderStages::all(),
+                    ref other => {
+                        log::error!("Intersection result of a pointer to {:?}", other);
+                        return Err(ExpressionError::InvalidRayQueryType(query));
+                    }
+                },
+                ref other => {
+                    log::error!("Intersection result of {:?}", other);
+                    return Err(ExpressionError::InvalidRayQueryType(query));
+                }
+            },
+            E::RayQueryVertexPositions {
+                query,
+                committed: _,
+            } => match resolver[query] {
+                Ti::Pointer {
+                    base,
+                    space: crate::AddressSpace::Function,
+                } => match resolver.types[base].inner {
+                    Ti::RayQuery {
+                        vertex_return: true,
+                    } => ShaderStages::all(),
                     ref other => {
                         log::error!("Intersection result of a pointer to {:?}", other);
                         return Err(ExpressionError::InvalidRayQueryType(query));

naga/src/valid/function.rs

@@ -173,6 +173,12 @@ pub enum FunctionError {
     InvalidRayQueryExpression(Handle<crate::Expression>),
     #[error("Acceleration structure {0:?} is not a matching expression")]
     InvalidAccelerationStructure(Handle<crate::Expression>),
+    #[error(
+        "Acceleration structure {0:?} is missing flag vertex_return while Ray Query {1:?} does"
+    )]
+    MissingAccelerationStructureVertexReturn(Handle<crate::Expression>, Handle<crate::Expression>),
+    #[error("Ray Query {0:?} is missing flag vertex_return")]
+    MissingRayQueryVertexReturn(Handle<crate::Expression>),
     #[error("Ray descriptor {0:?} is not a matching expression")]
     InvalidRayDescriptor(Handle<crate::Expression>),
     #[error("Ray Query {0:?} does not have a matching type")]
@@ -765,7 +771,8 @@ impl super::Validator {
                             | Ex::Math { .. }
                             | Ex::As { .. }
                             | Ex::ArrayLength(_)
-                            | Ex::RayQueryGetIntersection { .. } => {
+                            | Ex::RayQueryGetIntersection { .. }
+                            | Ex::RayQueryVertexPositions { .. } => {
                                 self.emit_expression(handle, context)?
                             }
                             Ex::CallResult(_)
@@ -1451,14 +1458,14 @@ impl super::Validator {
                                 .with_span_static(span, "invalid query expression"));
                         }
                     };
-                    match context.types[query_var.ty].inner {
-                        Ti::RayQuery => {}
+                    let rq_vertex_return = match context.types[query_var.ty].inner {
+                        Ti::RayQuery { vertex_return } => vertex_return,
                         ref other => {
                             log::error!("Unexpected ray query type {other:?}");
                             return Err(FunctionError::InvalidRayQueryType(query_var.ty)
                                 .with_span_static(span, "invalid query type"));
                         }
-                    }
+                    };
                     match *fun {
                         crate::RayQueryFunction::Initialize {
                             acceleration_structure,
@@ -1467,7 +1474,11 @@ impl super::Validator {
                             match *context
                                 .resolve_type(acceleration_structure, &self.valid_expression_set)?
                             {
-                                Ti::AccelerationStructure => {}
+                                Ti::AccelerationStructure { vertex_return } => {
+                                    if (!vertex_return) && rq_vertex_return {
+                                        return Err(FunctionError::MissingAccelerationStructureVertexReturn(acceleration_structure, query).with_span_static(span, "invalid acceleration structure"));
+                                    }
+                                }
                                 _ => {
                                     return Err(FunctionError::InvalidAccelerationStructure(
                                         acceleration_structure,

naga/src/valid/handles.rs

@@ -246,6 +246,9 @@ impl super::Validator {
         if let Some(ty) = special_types.ray_intersection {
             validate_type(ty)?;
         }
+        if let Some(ty) = special_types.ray_vertex_return {
+            validate_type(ty)?;
+        }
 
         for (handle, _node) in diagnostic_filters.iter() {
             let DiagnosticFilterNode { inner: _, parent } = diagnostic_filters[handle];
@@ -310,8 +313,8 @@ impl super::Validator {
             | crate::TypeInner::Atomic { .. }
             | crate::TypeInner::Image { .. }
             | crate::TypeInner::Sampler { .. }
-            | crate::TypeInner::AccelerationStructure
-            | crate::TypeInner::RayQuery => None,
+            | crate::TypeInner::AccelerationStructure { .. }
+            | crate::TypeInner::RayQuery { .. } => None,
             crate::TypeInner::Pointer { base, space: _ } => {
                 handle.check_dep(base)?;
                 None
@@ -559,6 +562,10 @@ impl super::Validator {
             crate::Expression::RayQueryGetIntersection {
                 query,
                 committed: _,
+            }
+            | crate::Expression::RayQueryVertexPositions {
+                query,
+                committed: _,
             } => {
                 handle.check_dep(query)?;
             }

naga/src/valid/interface.rs

@@ -551,8 +551,8 @@ impl super::Validator {
                         _ => {}
                     },
                     crate::TypeInner::Sampler { .. }
-                    | crate::TypeInner::AccelerationStructure
-                    | crate::TypeInner::RayQuery => {}
+                    | crate::TypeInner::AccelerationStructure { .. }
+                    | crate::TypeInner::RayQuery { .. } => {}
                     _ => {
                         return Err(GlobalVariableError::InvalidType(var.space));
                     }

naga/src/valid/mod.rs

@@ -161,6 +161,8 @@ bitflags::bitflags! {
         const TEXTURE_ATOMIC = 1 << 23;
         /// Support for atomic operations on 64-bit images.
         const TEXTURE_INT64_ATOMIC = 1 << 24;
+        /// Support for ray queries returning vertex position
+        const RAY_HIT_VERTEX_POSITION = 1 << 25;
     }
 }
 
@@ -403,8 +405,8 @@ impl crate::TypeInner {
             Self::Array { .. }
             | Self::Image { .. }
             | Self::Sampler { .. }
-            | Self::AccelerationStructure
-            | Self::RayQuery
+            | Self::AccelerationStructure { .. }
+            | Self::RayQuery { .. }
             | Self::BindingArray { .. } => false,
         }
     }

naga/src/valid/type.rs

@@ -704,15 +704,21 @@ impl super::Validator {
                 TypeFlags::ARGUMENT | TypeFlags::CREATION_RESOLVED,
                 Alignment::ONE,
             ),
-            Ti::AccelerationStructure => {
+            Ti::AccelerationStructure { vertex_return } => {
                 self.require_type_capability(Capabilities::RAY_QUERY)?;
+                if vertex_return {
+                    self.require_type_capability(Capabilities::RAY_HIT_VERTEX_POSITION)?;
+                }
                 TypeInfo::new(
                     TypeFlags::ARGUMENT | TypeFlags::CREATION_RESOLVED,
                     Alignment::ONE,
                 )
             }
-            Ti::RayQuery => {
+            Ti::RayQuery { vertex_return } => {
                 self.require_type_capability(Capabilities::RAY_QUERY)?;
+                if vertex_return {
+                    self.require_type_capability(Capabilities::RAY_HIT_VERTEX_POSITION)?;
+                }
                 TypeInfo::new(
                     TypeFlags::DATA
                         | TypeFlags::CONSTRUCTIBLE

naga/tests/out/ir/access.compact.ron

@@ -420,6 +420,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/access.ron

@@ -420,6 +420,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/collatz.compact.ron

@@ -43,6 +43,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/collatz.ron

@@ -43,6 +43,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/const_assert.compact.ron

@@ -3,6 +3,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/const_assert.ron

@@ -3,6 +3,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/diagnostic-filter.compact.ron

@@ -3,6 +3,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/diagnostic-filter.ron

@@ -3,6 +3,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/fetch_depth.compact.ron

@@ -66,6 +66,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [

naga/tests/out/ir/fetch_depth.ron

@@ -129,6 +129,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [

naga/tests/out/ir/index-by-value.compact.ron

@@ -80,6 +80,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/index-by-value.ron

@@ -80,6 +80,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/local-const.compact.ron

@@ -35,6 +35,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [

naga/tests/out/ir/local-const.ron

@@ -35,6 +35,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [

naga/tests/out/ir/must-use.compact.ron

@@ -11,6 +11,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/must-use.ron

@@ -11,6 +11,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron

@@ -52,6 +52,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {
             AtomicCompareExchangeWeakResult((
                 kind: Uint,

naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron

@@ -52,6 +52,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {
             AtomicCompareExchangeWeakResult((
                 kind: Uint,

naga/tests/out/ir/overrides-ray-query.compact.ron

@@ -9,11 +9,15 @@
         ),
         (
             name: None,
-            inner: AccelerationStructure,
+            inner: AccelerationStructure(
+                vertex_return: false,
+            ),
         ),
         (
             name: None,
-            inner: RayQuery,
+            inner: RayQuery(
+                vertex_return: false,
+            ),
         ),
         (
             name: None,
@@ -80,6 +84,7 @@
     special_types: (
         ray_desc: Some(5),
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/overrides-ray-query.ron

@@ -9,11 +9,15 @@
         ),
         (
             name: None,
-            inner: AccelerationStructure,
+            inner: AccelerationStructure(
+                vertex_return: false,
+            ),
         ),
         (
             name: None,
-            inner: RayQuery,
+            inner: RayQuery(
+                vertex_return: false,
+            ),
         ),
         (
             name: None,
@@ -80,6 +84,7 @@
     special_types: (
         ray_desc: Some(5),
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/overrides.compact.ron

@@ -25,6 +25,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/overrides.ron

@@ -25,6 +25,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/shadow.compact.ron

@@ -154,6 +154,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [

naga/tests/out/ir/shadow.ron

@@ -277,6 +277,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [

naga/tests/out/ir/spec-constants.compact.ron

@@ -170,6 +170,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [

naga/tests/out/ir/spec-constants.ron

@@ -261,6 +261,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [

naga/tests/out/ir/storage-textures.compact.ron

@@ -70,6 +70,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

naga/tests/out/ir/storage-textures.ron

@@ -70,6 +70,7 @@
     special_types: (
         ray_desc: None,
         ray_intersection: None,
+        ray_vertex_return: None,
         predeclared_types: {},
     ),
     constants: [],

tests/gpu-tests/ray_tracing/as_build.rs

@@ -19,7 +19,11 @@ static UNBUILT_BLAS: GpuTestConfiguration = GpuTestConfiguration::new()
     .run_sync(unbuilt_blas);
 
 fn unbuilt_blas(ctx: TestingContext) {
-    let as_ctx = AsBuildContext::new(&ctx);
+    let as_ctx = AsBuildContext::new(
+        &ctx,
+        AccelerationStructureFlags::empty(),
+        AccelerationStructureFlags::empty(),
+    );
 
     // Build the TLAS package with an unbuilt BLAS.
     let mut encoder = ctx
@@ -49,7 +53,11 @@ static OUT_OF_ORDER_AS_BUILD: GpuTestConfiguration = GpuTestConfiguration::new()
     .run_sync(out_of_order_as_build);
 
 fn out_of_order_as_build(ctx: TestingContext) {
-    let as_ctx = AsBuildContext::new(&ctx);
+    let as_ctx = AsBuildContext::new(
+        &ctx,
+        AccelerationStructureFlags::empty(),
+        AccelerationStructureFlags::empty(),
+    );
 
     //
     // Encode the TLAS build before the BLAS build, but submit them in the right order.
@@ -80,7 +88,11 @@ fn out_of_order_as_build(ctx: TestingContext) {
     // Create a clean `AsBuildContext`
     //
 
-    let as_ctx = AsBuildContext::new(&ctx);
+    let as_ctx = AsBuildContext::new(
+        &ctx,
+        AccelerationStructureFlags::empty(),
+        AccelerationStructureFlags::empty(),
+    );
 
     //
     // Encode the BLAS build before the TLAS build, but submit them in the wrong order.
@@ -131,7 +143,11 @@ fn out_of_order_as_build_use(ctx: TestingContext) {
     // Create a clean `AsBuildContext`
     //
 
-    let as_ctx = AsBuildContext::new(&ctx);
+    let as_ctx = AsBuildContext::new(
+        &ctx,
+        AccelerationStructureFlags::empty(),
+        AccelerationStructureFlags::empty(),
+    );
 
     //
     // Build in the right order, then rebuild the BLAS so the TLAS is invalid, then use the TLAS.
@@ -340,3 +356,177 @@ fn build_with_transform(ctx: TestingContext) {
     );
     ctx.queue.submit([encoder_build.finish()]);
 }
+
+#[gpu_test]
+static ONLY_BLAS_VERTEX_RETURN: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(
+        TestParameters::default()
+            .test_features_limits()
+            .features(
+                wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
+                    | wgpu::Features::EXPERIMENTAL_RAY_QUERY
+                    | wgpu::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN,
+            )
+            // https://github.com/gfx-rs/wgpu/issues/6727
+            .skip(FailureCase::backend_adapter(wgpu::Backends::VULKAN, "AMD")),
+    )
+    .run_sync(only_blas_vertex_return);
+
+fn only_blas_vertex_return(ctx: TestingContext) {
+    // Set up BLAS with TLAS
+    let as_ctx = AsBuildContext::new(
+        &ctx,
+        AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN,
+        AccelerationStructureFlags::empty(),
+    );
+
+    let mut encoder_blas = ctx
+        .device
+        .create_command_encoder(&CommandEncoderDescriptor {
+            label: Some("BLAS 1"),
+        });
+
+    encoder_blas.build_acceleration_structures([&as_ctx.blas_build_entry()], []);
+
+    let mut encoder_tlas = ctx
+        .device
+        .create_command_encoder(&CommandEncoderDescriptor {
+            label: Some("TLAS 1"),
+        });
+
+    encoder_tlas.build_acceleration_structures([], [&as_ctx.tlas_package]);
+
+    ctx.queue
+        .submit([encoder_blas.finish(), encoder_tlas.finish()]);
+
+    // Create a bind-group containing a TLAS with a bind-group layout that requires vertex return,
+    // because only the BLAS and not the TLAS has `AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN`
+    // this is invalid.
+    {
+        let bind_group_layout = ctx
+            .device
+            .create_bind_group_layout(&BindGroupLayoutDescriptor {
+                label: None,
+                entries: &[BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: BindingType::AccelerationStructure {
+                        vertex_return: true,
+                    },
+                    count: None,
+                }],
+            });
+        fail(
+            &ctx.device,
+            || {
+                let _ = ctx.device.create_bind_group(&BindGroupDescriptor {
+                    label: None,
+                    layout: &bind_group_layout,
+                    entries: &[BindGroupEntry {
+                        binding: 0,
+                        resource: BindingResource::AccelerationStructure(
+                            as_ctx.tlas_package.tlas(),
+                        ),
+                    }],
+                });
+            },
+            None,
+        );
+        // drop these
+    }
+
+    // We then use it with a shader that does not require vertex return which should succeed.
+    {
+        //
+        // Create shader to use tlas with
+        //
+
+        let shader = ctx
+            .device
+            .create_shader_module(include_wgsl!("shader.wgsl"));
+        let compute_pipeline = ctx
+            .device
+            .create_compute_pipeline(&ComputePipelineDescriptor {
+                label: None,
+                layout: None,
+                module: &shader,
+                entry_point: Some("basic_usage"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+
+        let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
+            label: None,
+            layout: &compute_pipeline.get_bind_group_layout(0),
+            entries: &[BindGroupEntry {
+                binding: 0,
+                resource: BindingResource::AccelerationStructure(as_ctx.tlas_package.tlas()),
+            }],
+        });
+
+        //
+        // Use TLAS
+        //
+
+        let mut encoder_compute = ctx
+            .device
+            .create_command_encoder(&CommandEncoderDescriptor::default());
+        {
+            let mut pass = encoder_compute.begin_compute_pass(&ComputePassDescriptor {
+                label: None,
+                timestamp_writes: None,
+            });
+            pass.set_pipeline(&compute_pipeline);
+            pass.set_bind_group(0, Some(&bind_group), &[]);
+            pass.dispatch_workgroups(1, 1, 1)
+        }
+
+        ctx.queue.submit(Some(encoder_compute.finish()));
+    }
+}
+
+#[gpu_test]
+static ONLY_TLAS_VERTEX_RETURN: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(
+        TestParameters::default()
+            .test_features_limits()
+            .features(
+                wgpu::Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE
+                    | wgpu::Features::EXPERIMENTAL_RAY_QUERY
+                    | wgpu::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN,
+            )
+            // https://github.com/gfx-rs/wgpu/issues/6727
+            .skip(FailureCase::backend_adapter(wgpu::Backends::VULKAN, "AMD")),
+    )
+    .run_sync(only_tlas_vertex_return);
+
+fn only_tlas_vertex_return(ctx: TestingContext) {
+    // Set up BLAS with TLAS
+    let as_ctx = AsBuildContext::new(
+        &ctx,
+        AccelerationStructureFlags::empty(),
+        AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN,
+    );
+
+    let mut encoder_blas = ctx
+        .device
+        .create_command_encoder(&CommandEncoderDescriptor {
+            label: Some("BLAS 1"),
+        });
+
+    encoder_blas.build_acceleration_structures([&as_ctx.blas_build_entry()], []);
+
+    let mut encoder_tlas = ctx
+        .device
+        .create_command_encoder(&CommandEncoderDescriptor {
+            label: Some("TLAS 1"),
+        });
+
+    fail(
+        &ctx.device,
+        || {
+            encoder_tlas.build_acceleration_structures([], [&as_ctx.tlas_package]);
+        },
+        None,
+    );
+}

tests/gpu-tests/ray_tracing/mod.rs

@@ -26,7 +26,11 @@ pub struct AsBuildContext {
 }
 
 impl AsBuildContext {
-    pub fn new(ctx: &TestingContext) -> Self {
+    pub fn new(
+        ctx: &TestingContext,
+        additional_blas_flags: AccelerationStructureFlags,
+        additional_tlas_flags: AccelerationStructureFlags,
+    ) -> Self {
         let vertices = ctx.device.create_buffer_init(&BufferInitDescriptor {
             label: None,
             contents: &[0; mem::size_of::<[[f32; 3]; 3]>()],
@@ -44,7 +48,7 @@ impl AsBuildContext {
         let blas = ctx.device.create_blas(
             &CreateBlasDescriptor {
                 label: Some("BLAS"),
-                flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
+                flags: AccelerationStructureFlags::PREFER_FAST_TRACE | additional_blas_flags,
                 update_mode: AccelerationStructureUpdateMode::Build,
             },
             BlasGeometrySizeDescriptors::Triangles {
@@ -55,7 +59,7 @@ impl AsBuildContext {
         let tlas = ctx.device.create_tlas(&CreateTlasDescriptor {
             label: Some("TLAS"),
             max_instances: 1,
-            flags: AccelerationStructureFlags::PREFER_FAST_TRACE,
+            flags: AccelerationStructureFlags::PREFER_FAST_TRACE | additional_tlas_flags,
             update_mode: AccelerationStructureUpdateMode::Build,
         });
 

tests/gpu-tests/ray_tracing/shader.rs

@@ -1,9 +1,9 @@
 use crate::ray_tracing::AsBuildContext;
-use wgpu::BufferUsages;
 use wgpu::{
     include_wgsl, BindGroupDescriptor, BindGroupEntry, BindingResource, BufferDescriptor,
     CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor,
 };
+use wgpu::{AccelerationStructureFlags, BufferUsages};
 use wgpu_macros::gpu_test;
 use wgpu_test::{GpuTestConfiguration, TestParameters, TestingContext};
 
@@ -28,7 +28,11 @@ fn access_all_struct_members(ctx: TestingContext) {
     // Create a clean `AsBuildContext`
     //
 
-    let as_ctx = AsBuildContext::new(&ctx);
+    let as_ctx = AsBuildContext::new(
+        &ctx,
+        AccelerationStructureFlags::empty(),
+        AccelerationStructureFlags::empty(),
+    );
 
     let mut encoder_build = ctx
         .device

wgpu-core/src/binding_model.rs

@@ -188,6 +188,8 @@ pub enum CreateBindGroupError {
         layout_flt: bool,
         sampler_flt: bool,
     },
+    #[error("TLAS binding {binding} is required to support vertex returns but is missing flag AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN")]
+    MissingTLASVertexReturn { binding: u32 },
     #[error("Bound texture views can not have both depth and stencil aspects enabled")]
     DepthStencilAspect,
     #[error("The adapter does not support read access for storage textures of format {0:?}")]
@@ -380,7 +382,7 @@ impl BindingTypeMaxCountValidator {
                 wgt::BindingType::StorageTexture { .. } => {
                     self.storage_textures.add(binding.visibility, count);
                 }
-                wgt::BindingType::AccelerationStructure => {
+                wgt::BindingType::AccelerationStructure { .. } => {
                     self.acceleration_structures.add(binding.visibility, count);
                 }
             }

wgpu-core/src/command/ray_tracing.rs

@@ -530,6 +530,21 @@ impl Global {
                     },
                 ));
 
+                if tlas
+                    .flags
+                    .contains(wgpu_types::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN)
+                    && !blas.flags.contains(
+                        wgpu_types::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN,
+                    )
+                {
+                    return Err(
+                        BuildAccelerationStructureError::TlasDependentMissingVertexReturn(
+                            tlas.error_ident(),
+                            blas.error_ident(),
+                        ),
+                    );
+                }
+
                 instance_count += 1;
 
                 dependencies.push(blas.clone());

wgpu-core/src/device/mod.rs

@@ -469,6 +469,10 @@ pub fn create_validator(
         Caps::SUBGROUP_VERTEX_STAGE,
         features.contains(wgt::Features::SUBGROUP_VERTEX),
     );
+    caps.set(
+        Caps::RAY_HIT_VERTEX_POSITION,
+        features.intersects(wgt::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN),
+    );
 
     naga::valid::Validator::new(flags, caps)
 }

wgpu-core/src/device/ray_tracing.rs

@@ -27,6 +27,13 @@ impl Device {
         self.check_is_valid()?;
         self.require_features(Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE)?;
 
+        if blas_desc
+            .flags
+            .contains(wgt::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN)
+        {
+            self.require_features(Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN)?;
+        }
+
         let size_info = match &sizes {
             wgt::BlasGeometrySizeDescriptors::Triangles { descriptors } => {
                 let mut entries =
@@ -139,6 +146,13 @@ impl Device {
             ));
         }
 
+        if desc
+            .flags
+            .contains(wgt::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN)
+        {
+            self.require_features(Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN)?;
+        }
+
         let size_info = unsafe {
             self.raw().get_acceleration_structure_build_sizes(
                 &hal::GetAccelerationStructureBuildSizesDescriptor {

wgpu-core/src/device/resource.rs

@@ -1905,7 +1905,7 @@ impl Device {
                         },
                     )
                 }
-                Bt::AccelerationStructure => (None, WritableStorage::No),
+                Bt::AccelerationStructure { .. } => (None, WritableStorage::No),
             };
 
             // Validate the count parameter
@@ -2249,7 +2249,15 @@ impl Device {
         tlas.same_device(self)?;
 
         match decl.ty {
-            wgt::BindingType::AccelerationStructure => (),
+            wgt::BindingType::AccelerationStructure { vertex_return } => {
+                if vertex_return
+                    && !tlas.flags.contains(
+                        wgpu_types::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN,
+                    )
+                {
+                    return Err(Error::MissingTLASVertexReturn { binding });
+                }
+            }
             _ => {
                 return Err(Error::WrongBindingType {
                     binding,

wgpu-core/src/ray_tracing.rs

@@ -130,6 +130,10 @@ pub enum BuildAccelerationStructureError {
 
     #[error("Blas {0:?} is missing the flag USE_TRANSFORM but the transform buffer is set")]
     UseTransformMissing(ResourceErrorIdent),
+    #[error(
+        "Tlas {0:?} dependent {1:?} is missing AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN"
+    )]
+    TlasDependentMissingVertexReturn(ResourceErrorIdent, ResourceErrorIdent),
 }
 
 #[derive(Clone, Debug, Error)]

wgpu-core/src/validation.rs

@@ -25,7 +25,9 @@ enum ResourceType {
     Sampler {
         comparison: bool,
     },
-    AccelerationStructure,
+    AccelerationStructure {
+        vertex_return: bool,
+    },
 }
 
 #[derive(Clone, Debug)]
@@ -42,7 +44,7 @@ impl From<&ResourceType> for BindingTypeName {
             ResourceType::Buffer { .. } => BindingTypeName::Buffer,
             ResourceType::Texture { .. } => BindingTypeName::Texture,
             ResourceType::Sampler { .. } => BindingTypeName::Sampler,
-            ResourceType::AccelerationStructure => BindingTypeName::AccelerationStructure,
+            ResourceType::AccelerationStructure { .. } => BindingTypeName::AccelerationStructure,
         }
     }
 }
@@ -54,7 +56,7 @@ impl From<&BindingType> for BindingTypeName {
             BindingType::Texture { .. } => BindingTypeName::Texture,
             BindingType::StorageTexture { .. } => BindingTypeName::Texture,
             BindingType::Sampler { .. } => BindingTypeName::Sampler,
-            BindingType::AccelerationStructure => BindingTypeName::AccelerationStructure,
+            BindingType::AccelerationStructure { .. } => BindingTypeName::AccelerationStructure,
         }
     }
 }
@@ -557,8 +559,10 @@ impl Resource {
                     });
                 }
             }
-            ResourceType::AccelerationStructure => match entry.ty {
-                BindingType::AccelerationStructure => (),
+            ResourceType::AccelerationStructure { vertex_return } => match entry.ty {
+                BindingType::AccelerationStructure {
+                    vertex_return: entry_vertex_return,
+                } if vertex_return == entry_vertex_return => (),
                 _ => {
                     return Err(BindingError::WrongType {
                         binding: (&entry.ty).into(),
@@ -650,7 +654,9 @@ impl Resource {
                     },
                 }
             }
-            ResourceType::AccelerationStructure => BindingType::AccelerationStructure,
+            ResourceType::AccelerationStructure { vertex_return } => {
+                BindingType::AccelerationStructure { vertex_return }
+            }
         })
     }
 }
@@ -949,7 +955,9 @@ impl Interface {
                     class,
                 },
                 naga::TypeInner::Sampler { comparison } => ResourceType::Sampler { comparison },
-                naga::TypeInner::AccelerationStructure => ResourceType::AccelerationStructure,
+                naga::TypeInner::AccelerationStructure { vertex_return } => {
+                    ResourceType::AccelerationStructure { vertex_return }
+                }
                 ref other => ResourceType::Buffer {
                     size: wgt::BufferSize::new(other.size(module.to_ctx()) as u64).unwrap(),
                 },

wgpu-hal/examples/ray-traced-triangle/main.rs

@@ -342,7 +342,9 @@ impl<A: hal::Api> Example<A> {
                 wgpu_types::BindGroupLayoutEntry {
                     binding: 2,
                     visibility: wgpu_types::ShaderStages::COMPUTE,
-                    ty: wgpu_types::BindingType::AccelerationStructure,
+                    ty: wgpu_types::BindingType::AccelerationStructure {
+                        vertex_return: false,
+                    },
                     count: None,
                 },
             ],

wgpu-hal/src/dx12/conv.rs

@@ -114,7 +114,7 @@ pub fn map_binding_type(ty: &wgt::BindingType) -> Direct3D12::D3D12_DESCRIPTOR_R
             ..
         }
         | Bt::StorageTexture { .. } => Direct3D12::D3D12_DESCRIPTOR_RANGE_TYPE_UAV,
-        Bt::AccelerationStructure => Direct3D12::D3D12_DESCRIPTOR_RANGE_TYPE_SRV,
+        Bt::AccelerationStructure { .. } => Direct3D12::D3D12_DESCRIPTOR_RANGE_TYPE_SRV,
     }
 }
 

wgpu-hal/src/dx12/device.rs

@@ -773,7 +773,7 @@ impl crate::Device for super::Device {
                 wgt::BindingType::Buffer { .. }
                 | wgt::BindingType::Texture { .. }
                 | wgt::BindingType::StorageTexture { .. }
-                | wgt::BindingType::AccelerationStructure => num_views += count,
+                | wgt::BindingType::AccelerationStructure { .. } => num_views += count,
                 wgt::BindingType::Sampler { .. } => has_sampler_in_group = true,
             }
         }
@@ -1514,7 +1514,7 @@ impl crate::Device for super::Device {
                         sampler_indexes.push(data.index);
                     }
                 }
-                wgt::BindingType::AccelerationStructure => {
+                wgt::BindingType::AccelerationStructure { .. } => {
                     let start = entry.resource_index as usize;
                     let end = start + entry.count as usize;
                     for data in &desc.acceleration_structures[start..end] {

wgpu-hal/src/gles/device.rs

@@ -1191,7 +1191,7 @@ impl crate::Device for super::Device {
                         ty: wgt::BufferBindingType::Storage { .. },
                         ..
                     } => &mut num_storage_buffers,
-                    wgt::BindingType::AccelerationStructure => unimplemented!(),
+                    wgt::BindingType::AccelerationStructure { .. } => unimplemented!(),
                 };
 
                 binding_to_slot[entry.binding as usize] = *counter;
@@ -1301,7 +1301,7 @@ impl crate::Device for super::Device {
                         format: format_desc.internal,
                     })
                 }
-                wgt::BindingType::AccelerationStructure => unimplemented!(),
+                wgt::BindingType::AccelerationStructure { .. } => unimplemented!(),
             };
             contents.push(binding);
         }

wgpu-hal/src/metal/device.rs

@@ -735,7 +735,7 @@ impl crate::Device for super::Device {
                                     wgt::StorageTextureAccess::Atomic => true,
                                 };
                             }
-                            wgt::BindingType::AccelerationStructure => unimplemented!(),
+                            wgt::BindingType::AccelerationStructure { .. } => unimplemented!(),
                         }
                     }
 
@@ -960,7 +960,7 @@ impl crate::Device for super::Device {
                                 );
                                 counter.textures += 1;
                             }
-                            wgt::BindingType::AccelerationStructure => unimplemented!(),
+                            wgt::BindingType::AccelerationStructure { .. } => unimplemented!(),
                         }
                     }
                 }

wgpu-hal/src/vulkan/adapter.rs

@@ -106,6 +106,7 @@ pub struct PhysicalDeviceFeatures {
     /// to Vulkan 1.3.
     zero_initialize_workgroup_memory:
         Option<vk::PhysicalDeviceZeroInitializeWorkgroupMemoryFeatures<'static>>,
+    position_fetch: Option<vk::PhysicalDeviceRayTracingPositionFetchFeaturesKHR<'static>>,
 
     /// Features provided by `VK_KHR_shader_atomic_int64`, promoted to Vulkan 1.2.
     shader_atomic_int64: Option<vk::PhysicalDeviceShaderAtomicInt64Features<'static>>,
@@ -173,6 +174,9 @@ impl PhysicalDeviceFeatures {
         if let Some(ref mut feature) = self.shader_atomic_int64 {
             info = info.push_next(feature);
         }
+        if let Some(ref mut feature) = self.position_fetch {
+            info = info.push_next(feature);
+        }
         if let Some(ref mut feature) = self.shader_image_atomic_int64 {
             info = info.push_next(feature);
         }
@@ -482,6 +486,14 @@ impl PhysicalDeviceFeatures {
             } else {
                 None
             },
+            position_fetch: if enabled_extensions.contains(&khr::ray_tracing_position_fetch::NAME) {
+                Some(
+                    vk::PhysicalDeviceRayTracingPositionFetchFeaturesKHR::default()
+                        .ray_tracing_position_fetch(true),
+                )
+            } else {
+                None
+            },
             mesh_shader: if enabled_extensions.contains(&ext::mesh_shader::NAME) {
                 let needed = requested_features.contains(wgt::Features::MESH_SHADER);
                 Some(
@@ -644,6 +656,10 @@ impl PhysicalDeviceFeatures {
             F::CONSERVATIVE_RASTERIZATION,
             caps.supports_extension(ext::conservative_rasterization::NAME),
         );
+        features.set(
+            F::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN,
+            caps.supports_extension(khr::ray_tracing_position_fetch::NAME),
+        );
 
         if let Some(ref descriptor_indexing) = self.descriptor_indexing {
             // We use update-after-bind descriptors for all bind groups containing binding arrays.
@@ -1049,6 +1065,10 @@ impl PhysicalDeviceProperties {
             extensions.push(khr::ray_query::NAME);
         }
 
+        if requested_features.contains(wgt::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN) {
+            extensions.push(khr::ray_tracing_position_fetch::NAME)
+        }
+
         // Require `VK_EXT_conservative_rasterization` if the associated feature was requested
         if requested_features.contains(wgt::Features::CONSERVATIVE_RASTERIZATION) {
             extensions.push(ext::conservative_rasterization::NAME);
@@ -1455,6 +1475,13 @@ impl super::InstanceShared {
                 features2 = features2.push_next(next);
             }
 
+            if capabilities.supports_extension(khr::ray_tracing_position_fetch::NAME) {
+                let next = features
+                    .position_fetch
+                    .insert(vk::PhysicalDeviceRayTracingPositionFetchFeaturesKHR::default());
+                features2 = features2.push_next(next);
+            }
+
             // `VK_KHR_zero_initialize_workgroup_memory` is promoted to 1.3
             if capabilities.device_api_version >= vk::API_VERSION_1_3
                 || capabilities.supports_extension(khr::zero_initialize_workgroup_memory::NAME)
@@ -1959,6 +1986,9 @@ impl super::Adapter {
             if features.contains(wgt::Features::EXPERIMENTAL_RAY_QUERY) {
                 capabilities.push(spv::Capability::RayQueryKHR);
             }
+            if features.contains(wgt::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN) {
+                capabilities.push(spv::Capability::RayQueryPositionFetchKHR)
+            }
             spv::Options {
                 lang_version: if features
                     .intersects(wgt::Features::SUBGROUP | wgt::Features::SUBGROUP_VERTEX)

wgpu-hal/src/vulkan/conv.rs

@@ -771,7 +771,9 @@ pub fn map_binding_type(ty: wgt::BindingType) -> vk::DescriptorType {
         wgt::BindingType::Sampler { .. } => vk::DescriptorType::SAMPLER,
         wgt::BindingType::Texture { .. } => vk::DescriptorType::SAMPLED_IMAGE,
         wgt::BindingType::StorageTexture { .. } => vk::DescriptorType::STORAGE_IMAGE,
-        wgt::BindingType::AccelerationStructure => vk::DescriptorType::ACCELERATION_STRUCTURE_KHR,
+        wgt::BindingType::AccelerationStructure { .. } => {
+            vk::DescriptorType::ACCELERATION_STRUCTURE_KHR
+        }
     }
 }
 
@@ -953,6 +955,10 @@ pub fn map_acceleration_structure_flags(
         vk_flags |= vk::BuildAccelerationStructureFlagsKHR::ALLOW_COMPACTION
     }
 
+    if flags.contains(crate::AccelerationStructureBuildFlags::ALLOW_RAY_HIT_VERTEX_RETURN) {
+        vk_flags |= vk::BuildAccelerationStructureFlagsKHR::ALLOW_DATA_ACCESS
+    }
+
     vk_flags
 }
 

wgpu-hal/src/vulkan/device.rs

@@ -1446,7 +1446,7 @@ impl crate::Device for super::Device {
                 wgt::BindingType::StorageTexture { .. } => {
                     desc_count.storage_image += count;
                 }
-                wgt::BindingType::AccelerationStructure => {
+                wgt::BindingType::AccelerationStructure { .. } => {
                     desc_count.acceleration_structure += count;
                 }
             }

wgpu-types/src/features.rs

@@ -1165,6 +1165,20 @@ bitflags_array! {
         ///
         /// This is a native only feature.
         const MESH_SHADER = 1 << 47;
+        /// ***THIS IS EXPERIMENTAL:*** Features enabled by this may have
+        /// major bugs in them and are expected to be subject to breaking changes, suggestions
+        /// for the API exposed by this should be posted on [the ray-tracing issue](https://github.com/gfx-rs/wgpu/issues/6762)
+        ///
+        /// Allows for returning of the hit triangle's vertex position when tracing with an
+        /// acceleration structure marked with [`AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN`].
+        ///
+        /// Supported platforms:
+        /// - Vulkan
+        ///
+        /// This is a native only feature
+        ///
+        /// [`AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN`]: super::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN
+        const EXPERIMENTAL_RAY_HIT_VERTEX_RETURN = 1 << 48;
     }
 
     /// Features that are not guaranteed to be supported.

wgpu-types/src/lib.rs

@@ -6495,12 +6495,24 @@ pub enum BindingType {
     /// var as: acceleration_structure;
     /// ```
     ///
+    /// or with vertex return enabled
+    /// ```rust,ignore
+    /// @group(0) @binding(0)
+    /// var as: acceleration_structure<vertex_return>;
+    /// ```
+    ///
     /// Example GLSL syntax:
     /// ```cpp,ignore
     /// layout(binding = 0)
     /// uniform accelerationStructureEXT as;
     /// ```
-    AccelerationStructure,
+    AccelerationStructure {
+        /// Whether this acceleration structure can be used to
+        /// create a ray query that has flag vertex return in the shader
+        ///
+        /// If enabled requires [`Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN`]
+        vertex_return: bool,
+    },
 }
 
 impl BindingType {
@@ -7283,6 +7295,8 @@ bitflags::bitflags!(
         /// Use `BlasTriangleGeometry::transform_buffer` when building a BLAS (only allowed in
         /// BLAS creation)
         const USE_TRANSFORM = 1 << 5;
+        /// Allow retrieval of the vertices of the triangle hit by a ray.
+        const ALLOW_RAY_HIT_VERTEX_RETURN = 1 << 6;
     }
 );
 

wgpu/src/backend/webgpu.rs

@@ -1961,7 +1961,7 @@ impl dispatch::DeviceInterface for WebDevice {
                             .set_view_dimension(map_texture_view_dimension(view_dimension));
                         mapped_entry.set_storage_texture(&storage_texture);
                     }
-                    wgt::BindingType::AccelerationStructure => todo!(),
+                    wgt::BindingType::AccelerationStructure { .. } => todo!(),
                 }
 
                 mapped_entry