diff --git a/CHANGELOG.md b/CHANGELOG.md index b3fb69137..14b664cfe 100644 --- a/CHANGELOG.md +++ b/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). diff --git a/docs/api-specs/ray_tracing.md b/docs/api-specs/ray_tracing.md index 54eae7540..9d56cf1a6 100644 --- a/docs/api-specs/ray_tracing.md +++ b/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, acceleration_structure: acceleration_structure, ray_desc: RayDesc) +// Overload. +rayQueryInitialize(rq: ptr>, acceleration_structure: acceleration_structure, 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, 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) -> bool +// Overload. +rayQueryProceed(rq: ptr>) -> 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) -> RayIntersection +// Overload. +rayQueryGetCommittedIntersection(rq: ptr>) -> RayIntersection // - Returns intersection details about a hit considered `Candidate`. rayQueryGetCandidateIntersection(rq: ptr) -> RayIntersection +// Overload. +rayQueryGetCandidateIntersection(rq: ptr>) -> RayIntersection + +// - Returns the vertices of the hit triangle considered `Committed`. +getCommittedHitVertexPositions(rq: ptr>) -> array, 3> + +// - Returns the vertices of the hit triangle considered `Candidate`. +getCandidateHitVertexPositions(rq: ptr>) -> array, 3> ``` > [!CAUTION] @@ -89,6 +103,11 @@ rayQueryGetCandidateIntersection(rq: ptr) -> 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`, diff --git a/examples/features/src/lib.rs b/examples/features/src/lib.rs index 6efca3668..f56f19c62 100644 --- a/examples/features/src/lib.rs +++ b/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; diff --git a/examples/features/src/main.rs b/examples/features/src/main.rs index a5b4ad673..d803ba249 100644 --- a/examples/features/src/main.rs +++ b/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 { diff --git a/examples/features/src/ray_cube_normals/README.md b/examples/features/src/ray_cube_normals/README.md new file mode 100644 index 000000000..3d33fc255 --- /dev/null +++ b/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) diff --git a/examples/features/src/ray_cube_normals/blit.wgsl b/examples/features/src/ray_cube_normals/blit.wgsl new file mode 100644 index 000000000..69adbb3cc --- /dev/null +++ b/examples/features/src/ray_cube_normals/blit.wgsl @@ -0,0 +1,52 @@ +struct VertexOutput { + @builtin(position) position: vec4, + @location(0) tex_coords: vec2, +}; + +// 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(x) * 2.0, + f32(y) * 2.0 + ); + result.position = vec4( + 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; +@group(0) +@binding(1) +var r_sampler: sampler; + +@fragment +fn fs_main(vertex: VertexOutput) -> @location(0) vec4 { + return textureSample(r_color, r_sampler, vertex.tex_coords); +} diff --git a/examples/features/src/ray_cube_normals/mod.rs b/examples/features/src/ray_cube_normals/mod.rs new file mode 100644 index 000000000..342afd374 --- /dev/null +++ b/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, Vec) { + 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` 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 { + inner: F, +} +impl>> Future for ErrorFuture { + 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::() 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::("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: ::required_features(), + required_limits: ::required_limits(), + force_fxc: false, + skips: vec![], + failures: Vec::new(), + required_downlevel_caps: + ::required_downlevel_capabilities(), + }, + comparisons: &[wgpu_test::ComparisonType::Mean(0.02)], + _phantom: std::marker::PhantomData::, +}; diff --git a/examples/features/src/ray_cube_normals/screenshot.png b/examples/features/src/ray_cube_normals/screenshot.png new file mode 100644 index 000000000..06d3fb1d1 Binary files /dev/null and b/examples/features/src/ray_cube_normals/screenshot.png differ diff --git a/examples/features/src/ray_cube_normals/shader.wgsl b/examples/features/src/ray_cube_normals/shader.wgsl new file mode 100644 index 000000000..93b69834e --- /dev/null +++ b/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, + dir: vec3, +} + +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, + front_face: bool, + object_to_world: mat4x3, + world_to_object: mat4x3, +} +*/ + +struct Uniforms { + view_inv: mat4x4, + proj_inv: mat4x4, +}; + +@group(0) @binding(0) +var output: texture_storage_2d; + +@group(0) @binding(1) +var uniforms: Uniforms; + +@group(0) @binding(2) +var acc_struct: acceleration_structure; + +@compute @workgroup_size(8, 8) +fn main(@builtin(global_invocation_id) global_id: vec3) { + let target_size = textureDimensions(output); + var color = vec4(vec2(global_id.xy) / vec2(target_size), 0.0, 1.0); + + + let pixel_center = vec2(global_id.xy) + vec2(0.5); + let in_uv = pixel_center/vec2(target_size.xy); + let d = in_uv * 2.0 - 1.0; + + let origin = (uniforms.view_inv * vec4(0.0,0.0,0.0,1.0)).xyz; + let temp = uniforms.proj_inv * vec4(d.x, d.y, 1.0, 1.0); + let direction = (uniforms.view_inv * vec4(normalize(temp.xyz), 0.0)).xyz; + + var rq: ray_query; + 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 = 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); +} diff --git a/examples/features/src/ray_shadows/mod.rs b/examples/features/src/ray_shadows/mod.rs index 944f31554..3e0a57f19 100644 --- a/examples/features/src/ray_shadows/mod.rs +++ b/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, }, ], diff --git a/examples/features/src/ray_traced_triangle/mod.rs b/examples/features/src/ray_traced_triangle/mod.rs index 695a7158c..99a82dd8e 100644 --- a/examples/features/src/ray_traced_triangle/mod.rs +++ b/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, }, ], diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs index 7a4104076..dd8246f90 100644 --- a/naga/src/back/dot/mod.rs +++ b/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 diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 4d434d70a..c669f032c 100644 --- a/naga/src/back/glsl/mod.rs +++ b/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(()) diff --git a/naga/src/back/hlsl/ray.rs b/naga/src/back/hlsl/ray.rs index 99cdaa64a..24eae4b57 100644 --- a/naga/src/back/hlsl/ray.rs +++ b/naga/src/back/hlsl/ray.rs @@ -28,7 +28,12 @@ impl 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 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())?; diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index be1bd8b3f..953def189 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/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")?; } @@ -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 { .. } diff --git a/naga/src/back/mod.rs b/naga/src/back/mod.rs index cead56d48..2acb18626 100644 --- a/naga/src/back/mod.rs +++ b/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, } } diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 5cf369541..3b1e28f67 100644 --- a/naga/src/back/msl/writer.rs +++ b/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 Writer { write!(self.out, ")")?; } } + crate::Expression::RayQueryVertexPositions { .. } => { + unimplemented!() + } crate::Expression::RayQueryGetIntersection { query, committed: _, @@ -3815,12 +3824,12 @@ impl Writer { 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); } diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index e601aca20..4a54da5b1 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -638,6 +638,12 @@ fn adjust_expr(new_pos: &HandleVec>, expr: &mut E | Expression::WorkGroupUniformLoadResult { ty: _ } | Expression::SubgroupBallotResult | Expression::SubgroupOperationResult { .. } => {} + Expression::RayQueryVertexPositions { + ref mut query, + committed: _, + } => { + adjust(query); + } } } diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 8c21f0533..006ab310e 100644 --- a/naga/src/back/spv/block.rs +++ b/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; diff --git a/naga/src/back/spv/instructions.rs b/naga/src/back/spv/instructions.rs index 749d9a312..db3934cb4 100644 --- a/naga/src/back/spv/instructions.rs +++ b/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, diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 985439691..bd12b2fbc 100644 --- a/naga/src/back/spv/mod.rs +++ b/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, diff --git a/naga/src/back/spv/ray.rs b/naga/src/back/spv/ray.rs index fb98b7081..41f50bf61 100644 --- a/naga/src/back/spv/ray.rs +++ b/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, + 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 + } } diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 9dadacc68..93756737d 100644 --- a/naga/src/back/spv/writer.rs +++ b/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); diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index b43dd9957..c12a6e6ee 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -628,7 +628,10 @@ impl Writer { } write!(self.out, ">")?; } - TypeInner::AccelerationStructure => write!(self.out, "acceleration_structure")?, + TypeInner::AccelerationStructure { vertex_return } => { + let caps = if vertex_return { "" } else { "" }; + write!(self.out, "acceleration_structure{}", caps)? + } _ => { return Err(Error::Unimplemented(format!("write_value_type {inner:?}"))); } @@ -1890,7 +1893,8 @@ impl Writer { 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 { .. } diff --git a/naga/src/compact/expressions.rs b/naga/src/compact/expressions.rs index 7ddb46619..d1b800a93 100644 --- a/naga/src/compact/expressions.rs +++ b/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), } } diff --git a/naga/src/compact/mod.rs b/naga/src/compact/mod.rs index 0fb2b0779..ee3d5d3f6 100644 --- a/naga/src/compact/mod.rs +++ b/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); } diff --git a/naga/src/compact/types.rs b/naga/src/compact/types.rs index cde1352d3..293256826 100644 --- a/naga/src/compact/types.rs +++ b/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 { diff --git a/naga/src/front/type_gen.rs b/naga/src/front/type_gen.rs index b13d271b3..687d245b1 100644 --- a/naga/src/front/type_gen.rs +++ b/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 { + 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 diff --git a/naga/src/front/wgsl/lower/conversion.rs b/naga/src/front/wgsl/lower/conversion.rs index 480fee5a9..58a7f3912 100644 --- a/naga/src/front/wgsl/lower/conversion.rs +++ b/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, } } diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 485f7c5b2..f40fe1547 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/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)) diff --git a/naga/src/front/wgsl/parse/ast.rs b/naga/src/front/wgsl/parse/ast.rs index f59354901..3318a3640 100644 --- a/naga/src/front/wgsl/parse/ast.rs +++ b/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 { diff --git a/naga/src/front/wgsl/parse/lexer.rs b/naga/src/front/wgsl/parse/lexer.rs index 955643e62..2c1f9608c 100644 --- a/naga/src/front/wgsl/parse/lexer.rs +++ b/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> { + 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('(')) } diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs index 501076026..8057a03ab 100644 --- a/naga/src/front/wgsl/parse/mod.rs +++ b/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), diff --git a/naga/src/front/wgsl/to_wgsl.rs b/naga/src/front/wgsl/to_wgsl.rs index acb79fd0f..f24e7296b 100644 --- a/naga/src/front/wgsl/to_wgsl.rs +++ b/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 { "" } else { "" }; + format!("acceleration_structure{}", caps) + } + Ti::RayQuery { vertex_return } => { + let caps = if 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"); diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 90f6fb99d..48156b977 100644 --- a/naga/src/lib.rs +++ b/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), + /// Get the Positions of the triangle hit by the [`RayQuery`] + /// + /// [`RayQuery`]: Statement::RayQuery + RayQueryVertexPositions { + query: Handle, + 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>, + /// Type for `RayVertexReturn + /// + /// Call [`Module::generate_vertex_return_type`] + pub ray_vertex_return: Option>, + /// Types for predeclared wgsl types instantiated on demand. /// /// Call [`Module::generate_predeclared_type`] to populate this if diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 6aaa98114..bb6893bc1 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/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), diff --git a/naga/src/proc/layouter.rs b/naga/src/proc/layouter.rs index 36d8883b8..0001d92f2 100644 --- a/naga/src/proc/layouter.rs +++ b/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, diff --git a/naga/src/proc/type_methods.rs b/naga/src/proc/type_methods.rs index d73b6df10..3b9e9348a 100644 --- a/naga/src/proc/type_methods.rs +++ b/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, } } diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs index f2764a5e4..2f189eb8b 100644 --- a/naga/src/proc/typifier.rs +++ b/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, diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs index e7732921c..5a3d6ebc3 100644 --- a/naga/src/valid/analyzer.rs +++ b/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))?; diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs index 32ebbb276..49b2fd27b 100644 --- a/naga/src/valid/expression.rs +++ b/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)); diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index b16aacffa..b6ff92bbd 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -173,6 +173,12 @@ pub enum FunctionError { InvalidRayQueryExpression(Handle), #[error("Acceleration structure {0:?} is not a matching expression")] InvalidAccelerationStructure(Handle), + #[error( + "Acceleration structure {0:?} is missing flag vertex_return while Ray Query {1:?} does" + )] + MissingAccelerationStructureVertexReturn(Handle, Handle), + #[error("Ray Query {0:?} is missing flag vertex_return")] + MissingRayQueryVertexReturn(Handle), #[error("Ray descriptor {0:?} is not a matching expression")] InvalidRayDescriptor(Handle), #[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, diff --git a/naga/src/valid/handles.rs b/naga/src/valid/handles.rs index 832fe2e76..93265e17a 100644 --- a/naga/src/valid/handles.rs +++ b/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)?; } diff --git a/naga/src/valid/interface.rs b/naga/src/valid/interface.rs index 5dcfac590..d182d9575 100644 --- a/naga/src/valid/interface.rs +++ b/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)); } diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs index ccb8709a2..e7a50d928 100644 --- a/naga/src/valid/mod.rs +++ b/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, } } diff --git a/naga/src/valid/type.rs b/naga/src/valid/type.rs index 404f8318f..5863eb813 100644 --- a/naga/src/valid/type.rs +++ b/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 diff --git a/naga/tests/out/ir/access.compact.ron b/naga/tests/out/ir/access.compact.ron index 53ba1a3c9..396376655 100644 --- a/naga/tests/out/ir/access.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/access.ron b/naga/tests/out/ir/access.ron index 53ba1a3c9..396376655 100644 --- a/naga/tests/out/ir/access.ron +++ b/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: [], diff --git a/naga/tests/out/ir/collatz.compact.ron b/naga/tests/out/ir/collatz.compact.ron index 6a7aebe54..6825677c7 100644 --- a/naga/tests/out/ir/collatz.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/collatz.ron b/naga/tests/out/ir/collatz.ron index 6a7aebe54..6825677c7 100644 --- a/naga/tests/out/ir/collatz.ron +++ b/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: [], diff --git a/naga/tests/out/ir/const_assert.compact.ron b/naga/tests/out/ir/const_assert.compact.ron index 15fa331fb..e15d22a02 100644 --- a/naga/tests/out/ir/const_assert.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/const_assert.ron b/naga/tests/out/ir/const_assert.ron index 15fa331fb..e15d22a02 100644 --- a/naga/tests/out/ir/const_assert.ron +++ b/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: [], diff --git a/naga/tests/out/ir/diagnostic-filter.compact.ron b/naga/tests/out/ir/diagnostic-filter.compact.ron index 315c8b342..18d1075ce 100644 --- a/naga/tests/out/ir/diagnostic-filter.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/diagnostic-filter.ron b/naga/tests/out/ir/diagnostic-filter.ron index 315c8b342..18d1075ce 100644 --- a/naga/tests/out/ir/diagnostic-filter.ron +++ b/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: [], diff --git a/naga/tests/out/ir/fetch_depth.compact.ron b/naga/tests/out/ir/fetch_depth.compact.ron index f10ccb94f..218ead0ef 100644 --- a/naga/tests/out/ir/fetch_depth.compact.ron +++ b/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: [ diff --git a/naga/tests/out/ir/fetch_depth.ron b/naga/tests/out/ir/fetch_depth.ron index dc9237beb..cc69b1c57 100644 --- a/naga/tests/out/ir/fetch_depth.ron +++ b/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: [ diff --git a/naga/tests/out/ir/index-by-value.compact.ron b/naga/tests/out/ir/index-by-value.compact.ron index 93a982142..691bbdcae 100644 --- a/naga/tests/out/ir/index-by-value.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/index-by-value.ron b/naga/tests/out/ir/index-by-value.ron index 93a982142..691bbdcae 100644 --- a/naga/tests/out/ir/index-by-value.ron +++ b/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: [], diff --git a/naga/tests/out/ir/local-const.compact.ron b/naga/tests/out/ir/local-const.compact.ron index 53fa73819..726874bb9 100644 --- a/naga/tests/out/ir/local-const.compact.ron +++ b/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: [ diff --git a/naga/tests/out/ir/local-const.ron b/naga/tests/out/ir/local-const.ron index 53fa73819..726874bb9 100644 --- a/naga/tests/out/ir/local-const.ron +++ b/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: [ diff --git a/naga/tests/out/ir/must-use.compact.ron b/naga/tests/out/ir/must-use.compact.ron index 3d51cb0c9..9024bd755 100644 --- a/naga/tests/out/ir/must-use.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/must-use.ron b/naga/tests/out/ir/must-use.ron index 3d51cb0c9..9024bd755 100644 --- a/naga/tests/out/ir/must-use.ron +++ b/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: [], diff --git a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron b/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron index 56be2f8ab..554558c8d 100644 --- a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron +++ b/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, diff --git a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron b/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron index 56be2f8ab..554558c8d 100644 --- a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron +++ b/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, diff --git a/naga/tests/out/ir/overrides-ray-query.compact.ron b/naga/tests/out/ir/overrides-ray-query.compact.ron index 10cad8353..f10fa8707 100644 --- a/naga/tests/out/ir/overrides-ray-query.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/overrides-ray-query.ron b/naga/tests/out/ir/overrides-ray-query.ron index 10cad8353..f10fa8707 100644 --- a/naga/tests/out/ir/overrides-ray-query.ron +++ b/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: [], diff --git a/naga/tests/out/ir/overrides.compact.ron b/naga/tests/out/ir/overrides.compact.ron index 00c57fa43..ddf2bb668 100644 --- a/naga/tests/out/ir/overrides.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/overrides.ron b/naga/tests/out/ir/overrides.ron index 00c57fa43..ddf2bb668 100644 --- a/naga/tests/out/ir/overrides.ron +++ b/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: [], diff --git a/naga/tests/out/ir/shadow.compact.ron b/naga/tests/out/ir/shadow.compact.ron index 24ab5edda..a490d4a48 100644 --- a/naga/tests/out/ir/shadow.compact.ron +++ b/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: [ diff --git a/naga/tests/out/ir/shadow.ron b/naga/tests/out/ir/shadow.ron index f7b4b67dc..24f973c82 100644 --- a/naga/tests/out/ir/shadow.ron +++ b/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: [ diff --git a/naga/tests/out/ir/spec-constants.compact.ron b/naga/tests/out/ir/spec-constants.compact.ron index d8658a413..5d53641ab 100644 --- a/naga/tests/out/ir/spec-constants.compact.ron +++ b/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: [ diff --git a/naga/tests/out/ir/spec-constants.ron b/naga/tests/out/ir/spec-constants.ron index 407ce5f49..ebca0aa83 100644 --- a/naga/tests/out/ir/spec-constants.ron +++ b/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: [ diff --git a/naga/tests/out/ir/storage-textures.compact.ron b/naga/tests/out/ir/storage-textures.compact.ron index 3f2f06439..ad04c885f 100644 --- a/naga/tests/out/ir/storage-textures.compact.ron +++ b/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: [], diff --git a/naga/tests/out/ir/storage-textures.ron b/naga/tests/out/ir/storage-textures.ron index 3f2f06439..ad04c885f 100644 --- a/naga/tests/out/ir/storage-textures.ron +++ b/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: [], diff --git a/tests/gpu-tests/ray_tracing/as_build.rs b/tests/gpu-tests/ray_tracing/as_build.rs index 4dbbd32e8..682740115 100644 --- a/tests/gpu-tests/ray_tracing/as_build.rs +++ b/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, + ); +} diff --git a/tests/gpu-tests/ray_tracing/mod.rs b/tests/gpu-tests/ray_tracing/mod.rs index a502ddccb..10b087496 100644 --- a/tests/gpu-tests/ray_tracing/mod.rs +++ b/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, }); diff --git a/tests/gpu-tests/ray_tracing/shader.rs b/tests/gpu-tests/ray_tracing/shader.rs index 08e415523..851cdc7f3 100644 --- a/tests/gpu-tests/ray_tracing/shader.rs +++ b/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 diff --git a/wgpu-core/src/binding_model.rs b/wgpu-core/src/binding_model.rs index 111442270..5f6d52e00 100644 --- a/wgpu-core/src/binding_model.rs +++ b/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); } } diff --git a/wgpu-core/src/command/ray_tracing.rs b/wgpu-core/src/command/ray_tracing.rs index 868d43ab1..f6334d827 100644 --- a/wgpu-core/src/command/ray_tracing.rs +++ b/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()); diff --git a/wgpu-core/src/device/mod.rs b/wgpu-core/src/device/mod.rs index 91f466d23..30c4daad6 100644 --- a/wgpu-core/src/device/mod.rs +++ b/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) } diff --git a/wgpu-core/src/device/ray_tracing.rs b/wgpu-core/src/device/ray_tracing.rs index 52d01b6a5..014674125 100644 --- a/wgpu-core/src/device/ray_tracing.rs +++ b/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 { diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs index 69a62ab11..b61bf1442 100644 --- a/wgpu-core/src/device/resource.rs +++ b/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, diff --git a/wgpu-core/src/ray_tracing.rs b/wgpu-core/src/ray_tracing.rs index 2d8217dfe..7810854ee 100644 --- a/wgpu-core/src/ray_tracing.rs +++ b/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)] diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs index 274cfab1c..175c47a06 100644 --- a/wgpu-core/src/validation.rs +++ b/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(), }, diff --git a/wgpu-hal/examples/ray-traced-triangle/main.rs b/wgpu-hal/examples/ray-traced-triangle/main.rs index e0083a309..9e2a7771e 100644 --- a/wgpu-hal/examples/ray-traced-triangle/main.rs +++ b/wgpu-hal/examples/ray-traced-triangle/main.rs @@ -342,7 +342,9 @@ impl Example { 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, }, ], diff --git a/wgpu-hal/src/dx12/conv.rs b/wgpu-hal/src/dx12/conv.rs index 3342a6ab0..b636dbc8d 100644 --- a/wgpu-hal/src/dx12/conv.rs +++ b/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, } } diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs index 53fc01cd9..87e611a18 100644 --- a/wgpu-hal/src/dx12/device.rs +++ b/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] { diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs index 76cb429ac..123d87f8f 100644 --- a/wgpu-hal/src/gles/device.rs +++ b/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); } diff --git a/wgpu-hal/src/metal/device.rs b/wgpu-hal/src/metal/device.rs index 49e7cb267..6fb172d00 100644 --- a/wgpu-hal/src/metal/device.rs +++ b/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!(), } } } diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs index 196eb85e7..e751464c2 100644 --- a/wgpu-hal/src/vulkan/adapter.rs +++ b/wgpu-hal/src/vulkan/adapter.rs @@ -106,6 +106,7 @@ pub struct PhysicalDeviceFeatures { /// to Vulkan 1.3. zero_initialize_workgroup_memory: Option>, + position_fetch: Option>, /// Features provided by `VK_KHR_shader_atomic_int64`, promoted to Vulkan 1.2. shader_atomic_int64: Option>, @@ -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) diff --git a/wgpu-hal/src/vulkan/conv.rs b/wgpu-hal/src/vulkan/conv.rs index 204d20ada..e1d5cb30e 100644 --- a/wgpu-hal/src/vulkan/conv.rs +++ b/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 } diff --git a/wgpu-hal/src/vulkan/device.rs b/wgpu-hal/src/vulkan/device.rs index 59d974177..d12748fcf 100644 --- a/wgpu-hal/src/vulkan/device.rs +++ b/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; } } diff --git a/wgpu-types/src/features.rs b/wgpu-types/src/features.rs index 229b845c7..1ffe3a016 100644 --- a/wgpu-types/src/features.rs +++ b/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. diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index f1f522637..e88773348 100644 --- a/wgpu-types/src/lib.rs +++ b/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; + /// ``` + /// /// 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; } ); diff --git a/wgpu/src/backend/webgpu.rs b/wgpu/src/backend/webgpu.rs index 508a72fbc..e2a26146d 100644 --- a/wgpu/src/backend/webgpu.rs +++ b/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