From 5b3266db23e8234d03e4d507b873f50ca4ab53ea Mon Sep 17 00:00:00 2001 From: Vecvec <130132884+Vecvec@users.noreply.github.com> Date: Wed, 5 Mar 2025 08:06:44 +1300 Subject: [PATCH] Support getting hit vertex positions (#7183) --- CHANGELOG.md | 2 + docs/api-specs/ray_tracing.md | 19 + examples/features/src/lib.rs | 1 + examples/features/src/main.rs | 6 + .../features/src/ray_cube_normals/README.md | 14 + .../features/src/ray_cube_normals/blit.wgsl | 52 ++ examples/features/src/ray_cube_normals/mod.rs | 492 ++++++++++++++++++ .../src/ray_cube_normals/screenshot.png | Bin 0 -> 30010 bytes .../features/src/ray_cube_normals/shader.wgsl | 85 +++ examples/features/src/ray_shadows/mod.rs | 4 +- .../features/src/ray_traced_triangle/mod.rs | 4 +- naga/src/back/dot/mod.rs | 5 + naga/src/back/glsl/mod.rs | 7 +- naga/src/back/hlsl/ray.rs | 14 +- naga/src/back/hlsl/writer.rs | 8 +- naga/src/back/mod.rs | 2 +- naga/src/back/msl/writer.rs | 21 +- naga/src/back/pipeline_constants.rs | 6 + naga/src/back/spv/block.rs | 7 + naga/src/back/spv/instructions.rs | 14 + naga/src/back/spv/mod.rs | 4 +- naga/src/back/spv/ray.rs | 48 +- naga/src/back/spv/writer.rs | 44 +- naga/src/back/wgsl/writer.rs | 8 +- naga/src/compact/expressions.rs | 10 + naga/src/compact/mod.rs | 9 + naga/src/compact/types.rs | 8 +- naga/src/front/type_gen.rs | 30 ++ naga/src/front/wgsl/lower/conversion.rs | 8 +- naga/src/front/wgsl/lower/mod.rs | 32 +- naga/src/front/wgsl/parse/ast.rs | 8 +- naga/src/front/wgsl/parse/lexer.rs | 22 + naga/src/front/wgsl/parse/mod.rs | 10 +- naga/src/front/wgsl/to_wgsl.rs | 10 +- naga/src/lib.rs | 17 +- naga/src/proc/constant_evaluator.rs | 4 +- naga/src/proc/layouter.rs | 4 +- naga/src/proc/type_methods.rs | 12 +- naga/src/proc/typifier.rs | 8 + naga/src/valid/analyzer.rs | 7 + naga/src/valid/expression.rs | 23 +- naga/src/valid/function.rs | 21 +- naga/src/valid/handles.rs | 11 +- naga/src/valid/interface.rs | 4 +- naga/src/valid/mod.rs | 6 +- naga/src/valid/type.rs | 10 +- naga/tests/out/ir/access.compact.ron | 1 + naga/tests/out/ir/access.ron | 1 + naga/tests/out/ir/collatz.compact.ron | 1 + naga/tests/out/ir/collatz.ron | 1 + naga/tests/out/ir/const_assert.compact.ron | 1 + naga/tests/out/ir/const_assert.ron | 1 + .../out/ir/diagnostic-filter.compact.ron | 1 + naga/tests/out/ir/diagnostic-filter.ron | 1 + naga/tests/out/ir/fetch_depth.compact.ron | 1 + naga/tests/out/ir/fetch_depth.ron | 1 + naga/tests/out/ir/index-by-value.compact.ron | 1 + naga/tests/out/ir/index-by-value.ron | 1 + naga/tests/out/ir/local-const.compact.ron | 1 + naga/tests/out/ir/local-const.ron | 1 + naga/tests/out/ir/must-use.compact.ron | 1 + naga/tests/out/ir/must-use.ron | 1 + ...ides-atomicCompareExchangeWeak.compact.ron | 1 + .../overrides-atomicCompareExchangeWeak.ron | 1 + .../out/ir/overrides-ray-query.compact.ron | 9 +- naga/tests/out/ir/overrides-ray-query.ron | 9 +- naga/tests/out/ir/overrides.compact.ron | 1 + naga/tests/out/ir/overrides.ron | 1 + naga/tests/out/ir/shadow.compact.ron | 1 + naga/tests/out/ir/shadow.ron | 1 + naga/tests/out/ir/spec-constants.compact.ron | 1 + naga/tests/out/ir/spec-constants.ron | 1 + .../tests/out/ir/storage-textures.compact.ron | 1 + naga/tests/out/ir/storage-textures.ron | 1 + tests/gpu-tests/ray_tracing/as_build.rs | 198 ++++++- tests/gpu-tests/ray_tracing/mod.rs | 10 +- tests/gpu-tests/ray_tracing/shader.rs | 8 +- wgpu-core/src/binding_model.rs | 4 +- wgpu-core/src/command/ray_tracing.rs | 15 + wgpu-core/src/device/mod.rs | 4 + wgpu-core/src/device/ray_tracing.rs | 14 + wgpu-core/src/device/resource.rs | 12 +- wgpu-core/src/ray_tracing.rs | 4 + wgpu-core/src/validation.rs | 22 +- wgpu-hal/examples/ray-traced-triangle/main.rs | 4 +- wgpu-hal/src/dx12/conv.rs | 2 +- wgpu-hal/src/dx12/device.rs | 4 +- wgpu-hal/src/gles/device.rs | 4 +- wgpu-hal/src/metal/device.rs | 4 +- wgpu-hal/src/vulkan/adapter.rs | 30 ++ wgpu-hal/src/vulkan/conv.rs | 8 +- wgpu-hal/src/vulkan/device.rs | 2 +- wgpu-types/src/features.rs | 14 + wgpu-types/src/lib.rs | 16 +- wgpu/src/backend/webgpu.rs | 2 +- 95 files changed, 1451 insertions(+), 115 deletions(-) create mode 100644 examples/features/src/ray_cube_normals/README.md create mode 100644 examples/features/src/ray_cube_normals/blit.wgsl create mode 100644 examples/features/src/ray_cube_normals/mod.rs create mode 100644 examples/features/src/ray_cube_normals/screenshot.png create mode 100644 examples/features/src/ray_cube_normals/shader.wgsl 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 0000000000000000000000000000000000000000..06d3fb1d19628f5a233113fa9e13ec43d0ba4b10 GIT binary patch literal 30010 zcmbV#c|a4__WytdT0ro%Dvcs&Tg5#JA|Ojpt7u)o6<31FqD3HrvIPh>PpK9KADk>l(0a*e92^b(imie7YfPM?U-+S?|33JCgcg{Wc-1GUIbFSwc zw%SeBoS}&z$mES1e%y{AU%<~v$d~GHdGci1WCS^-vhl~&I|JTzv;7l2-S;jXwy^(~ z?i{Tt)2_rU-ZcMM=lsMU_PWF;efL!p$7AOP)nqpLhspFQYw#W7sU3B9175F17R~+j zi));|$n0Yak6BKg_GZU|0zrR7&xOdAkmvf}H6sj>6rV`_*AJevsJ*H8Sjqp6X&5V--b;|IE3cEK$xd8`^1 z1J72$vz`?8H0PesZXPjVAknbpl0+p{N68&vxpKp?+O3@M4V9E;xN?&f5T~4U?B*-A!H{k#*}ak* z#slTH$7bb(W3DiZ|FdT`O*xw>i&suKsFRYaqvVcZr<`;Kmn-+V2}*i;5z6iEJwGZp zoXbi|XOD8-C)`p_HBTsMM)oSF3+N^#pWm4fD+!C5ajzJV2I(%uqvkKLdBvM56 z$(OPDqD9DN+Ip43JnU1g_2bg8$tJB=D@j88R^oh7r*&eQp^{`#h!**qRcr@07LXyPY%J%U0s&G(@0rXudm9&yO z8;ut=!=gb~K_xSlT~{oHswb&7I>G_~<4~upQky7j6c`7ZVr8CHR;exIn~eKTM53!$ z7MGRnVe_a`%jjSZ-Te~oDY+~H(nco}-6FRRG-Vf7K*DGbz*J&_4E)~&d2$(_I)c9h@yL8IcEst*U>l z!kb2SWzuzQ(F{yS@LM#)GF7L^uNVEijl+94MXAmnw>+?ip8BQVEZY9@;>>AJ{IBWn zcDfImr}}B-2PZ6V!>2@qgIN;2{f`nV8>~@5Lga#IrPSi7braJw%#hFt=o}&a|iiV zVZ2xyHN!TGyuPEJX}O8l%OD@LSey_wVJcw7DX%D-bXun&yisKZ6N0(oOOQ&?Lp@Lj z?5a@Kyon3OcjBV;1y@p8!!fU#rP^pUYP+27j@XUBjbJ;$ZopMxyKL|i>}P<}1QYIq zZJAo1Hj;_|-)$r=ebg1#l5e&-h&f>a;xy&1{A&EJpj7`k&uGG;VA)StvXzeili8rf z^K^40z<%91F;`JJ4Mn#KPielc_-wWueT^w$y^^W41P-Uwwo_e+D2+t}wExPz_ z2CH!OWEWDv_Mcnx(X~G>oPEdF{7ItJOh|(58H>b&56x2GA+2|K6Zx z|4*b>T%8w{mB+eSITKQi$V}ho#V<}+*NMNM_S3DTphx0lxBRy70T%(jJRm~AjVkc& z+|~q30N|z^TvyDc`b&h~{^*BZ?#CEPTP8k6s=Lx=DN}k@&wGfwx4>VN71M?GEu_2W zqPyu4qAu$gBS$v9P^t$GjFo^7?xFc9cFh3caX39a2gSCmHxK>i)P&9*%C{$?`qujS z2S&94Y_#0C^?baA1njag-Gr9z?mP2k7?u@q(X8$zT@s-({VO(ry z_$-`PGe(Y_A2ku+5AMKhC0rG*bLYgRfi;es4kvIwCrV0B0_G)VbZg&8XikcGgTeWpq z$5DY*UfuiI6;18s1L=q5%U9uU-du`gchc%w<_0F|NoP_r88rT>ZW`zLoH-Rul?L6+ zm4w}Ms3Aeui;iRm?un~B^qM+}`W5If{TOW;G#Dfw`gRe`8{AjHQMd=!7a6ZsjT>bpLa}&tDF||M!eouKXi? zzo7LK)jKQ8I9058T0s#by}g5B%9{^ z?t&K5iNlvHnW>S*&z76=B{H)Hk*GCS_%Bm_@Uxj8x&EFyX(xCdS^;Of9fjOvuMa0y zzxuaajoMU0oGS1RyV@GA6Vq7b&b*qFJw72ntCN|(w@wmn#C1t7>*z$3xMhem_RB6B z#uQc3SDD)yTq)wz~b;iPuhSD^iRR-$qVm0Qe2|}9IWeO}~ zs<_LGGpc=${!XWaUhXVdDb{-;^7M7` zh92D^OWD797TS8H$T_%S_R|B#`GXwV<5|oa7JhK{vQ7?Nz$Ooy^pxAu)A@OAQBsvtSo zo-wfI@n6B4RHhsiiKim5-A?VYncbl@)!8rj&G%Tt@2mn-3bfcE1yaEdM6l0*W>Hfdd?FB?971D2Vo8PckrHAGou<_>RdH7ASv--?#9@N`HnUwgHAJawX z%mn)pl`h#<~g%51IF^voz@#%knuY$G>+eClN3Ac4f;dX#CnXlz4>xA5Sb3hRUj4J_^Gai z4VR`N?8`tbR_z305y-tnNr$^i&pf&Iq3adYjk>K*e^Lsp_olzYX<2Pg(4Vn2=itx zLY~aRgntj1RhS0bp%KQtoV`b-sv8r=mR4-ufPtljPDhxEC^=}2)e;a2d52nIHSD@f z@W>QGk^ho2ncRs!nQS6)^pOnufD!Nf=bB+&cJRqZ7iz$OY*id&7@W(2u&&u|h{V3vlx@HVxzvT40Oghmi5`4lb8gT`b zpC1JU1oJ~WoT1vNdq*?un@Q6U3IsnCLXrQH3x%*@!t#;H10rG9H9IaQ5#Q?&bX)z);JDRcteMnH~)#{ zY6rlj05VEGu^UJiBOG}$hI}V-f#uT;HpN~Y3@va{BZhX&r?3Bnp`GZ{V;#y>mfx`M zt13SglIuqw^U}t^Sun9_p%}d{Vp}tedDh&6Dx`;U=i|UH5J6_>-Cm^|5>0H7)EKg?7%|a* zuJXyvJ|nfaX5>Y;1RCcKz4r?pKJ~u0!tm?+EhK*}v7v7;<4i3>&y0cuHS`8vKdf;e zIJlW5wVOrrp5qswB_{L@25)+pw7(XsK)3`sF|9)l7EYJjEcy=moiJc>dM!-B9Vu

Vq?+R~$FR^a+cxQ0Dd z_Vk49y$%_Bb^NCfR5OfLL>Aeox%Sjr=$`FtUAC~xgyb*@ryqL4f*45IFiEXrBu4Hp zPlqZG>8~R^lx?b7j~_Oq*bNcb8>A0twS((O1^Q`0c6&8s-GlCK<&9_kuZsKMUXoN3 z9uA!bfK=(<*X&ApM_6@4UEH^fYKQ*$R-?E>f4z0^yio2%oc`Nl_xDuF;p|6}UhANF ziiKMg{I&pZ6BT_EtuBtP?;dm<$d{zH_S6|L8oX*y?rZm^vjzqO+0C@L;l&OAY7f03 z3n3U^E0R)ezIofGk?d2;44+%;C$Zjv)1TfK(^GiQr!m`KzlJWAW0SuN(88$%CDBd1 z9IgD%^?QC0>O@`5bi5sNKMW298*%BqZ{!cvc|VW$ z)^_;y@gapXXlLX8T!4tL2F^-rce7qR+GLOs-P6YEE;O8-y{W0|KZuVU|)RykD#_&8z4$)MlK z%!g-H+mEU7WtB5@O7aXOX0f_8B1+wLO)ikLFx!!r^cAhQGh!Sw6S;f=I?Y z7_H9~5eS4T{wSi$Vw3Ylef$)1+K>yHtWg(Y)xVb*YosxlZ|_IMk`WTwddX?xcd~Hx zz{(|9N=Bjw<4Q0mIh}G@ArKARRHS37?A3~pc3Q$ItuTb1z90FbE`+T*yO2L~I-=so z%|O1)K+me#vi#*g$&iphSD?a5%`hX+=DEi_jVasSwHMMV)BNqHA*cPgQvNNZme04Y zrS!8ms#GPh&FU!iV}ra-HK(ux%TDqL1e%@B(Q z{%Jz|PinTh1CVGmKN!u^4EsftejAy85`71IsXzB;?D}*dy#SdnH;ItPfdRf#1wWWd zw(smOq;2(7vn>;|Vpn&o*{&9wK~RCDyJ8=5G6P+VgR2XJ(hTuWk*$ZCZMfL`>T38l z+L#|v(nkLFerFZInl3vD&R%wr-k;?ITxl}UlgDuyf2=azBdrFOF#ExzZEv)nQ*KU0 zO8iK(@6Vww6yvdMloYl{E#jQ?4y@%n#INWe2!BM)hIvIw9bJ>BMj9k4NJ_6^m7kM1HdO*#ViW|w2$eSr>@0nfu z9E7x0E9{iFYX}ijhT@LRjRhD@6|&l^BjIusd7+XHw}q($yk}$;&5EWq1YWu zY+)(vBFy2DCxIKExRC6OD4c_omAerAY6VD@WrCmNEXs!O#C|CV$h{N2IF_7%G&1q{ zKN(O%CZqU!BJP40H|+nUUxfwES%2#mz*JfAeIx)2lCZ!(QNUE}m%@+ZP4u1!6MYsL z7PFI!dIIh;gh66iC7(b@eVIrAh8+Rk0yu0g(hL)Ao`w(%AaG?5B2Zv=qfzWG_(aeI z8pmUItV3YQKjp1~;{SiJyOY9P19txk8wmzev7ikGNoxyIwRdv?iz_lrMBwj^%W}aQ z$Hfft9Y0qliketVCP%Q50Q4YC2bmDZWP&9%$OB1&i71r;jpJp4b@*8_2_2J(9d#Na zfOQ+X0`Qn@RB1yfFeW(&;?gDt7@=`|fD!9ZP8lhyvWgNwUw5^_2tbB5%Ya$g04h5P z)Kh_%PmeAi%It}73>wF$%wZjV7IN1ZTRarQ2gqO}fFJ?jWkaBr6wpH#)I$M1s1p|t z);JD&7>)#&l!hKQih?7I1fv5+4F}{-(SuNi!>FQ!V#rjkD25Pc9Y=g&jpGz$2rHfc zEJYddbuv|*=K>sSx(8R#lcrHUk&;IBga8E8xInZ78prDi*57C8$@g+c64R5vVC_VJ zJ+aeRh}c{iMPr6wMHCShJHXH5GU>3!akQ5W6mEPLfe@@HoaX~h4hCG91h}vnQxiyz zQm9G9n3^O_R1>Umyqb(eVtsSG=wBYqL~?yDL?`i14gOJkQimQ@Z2JW5uwJ@)chn@( z6N=Y5+jxtHDE+gqq!FFRzIv4L%oKn7%ez;KPwzhQ%gw8m*PXqt*lpCo`A752|d@Au-Vqk(UN)bqbG$J|bhp*#W#Vo10m{*NV()9ke@(U$Vmse`*TMs5St%W@ zchjhXRznX_X9VlsP%DX}pXPZ2-7zz0d4dr&c7FrVYjMDHrGqi17%I>x@TfrNw)LEgg=tjboFc= zK+`TW=A1e?Pt0m=YQ!u2g;yc`ZJ7ny6+!LvKKw`YNi)u_LbfKC4Vcpzjf%KGuqL0a z?KONMpHk4dBjTpOHe;8Lzbga1Nvm>(yf2cjq-t9Qz-AX2E?{t5|E%*hWsNX-jE4Xr{Q&_fdJN`;G;E zJmOXHDXwA40&0k6W#e+wsnn!jM5yd0;APGEVJkoS>+fRav|^H91U?7Th8!PFyX4eI zaOyejZ9?#%>J|3;V*d#8b9X_ARshP|#y zK;K2Rj=SNkKLxSNsl`Eqn}d7cMl-o@v%Zc@QR8>{edcA!^G&S2GB!F?CM#(0Y$IBH zjKzmW>D!2-rItNjZRnG&>K?wu#DPrY-3vnARvrH@@(kxO1$J*s8)ZcaUl9UYatzdK zd};}`CQam;QfE&ct%1z}jjJrI815o0XZrtpJPqXa!tpe1%a~s@p5RJC@Zx}r@w6;= zp5qsFA{==@nmSu-bBe|A{_2t78&Q`y!9ujel~>YTCu%l~=#2yLX)=5tp>>3B=@z%n z{u_PrGY!pZeVpEAMt8br8C|E`+4>>lk9sW|NTOlCnCCvRVhFfKs8Kj;ym{uOZ%-XR?q_Uzxu6@f7pC z_m#|pBo&`E&*Hhe+nYLWSlHAE>RWoQ3M=h)HVl3IMfZh3Po(P$vGW;ufHwsqK+=>*fCyK;Hi1|>g4QPAzInuI{fE!e};J!OZ0FF zUvl~)zvKJ;4dR!Z`5kNbH4Ht0s}!!+%eC;ZY4}5(S9GgzU+}qiM=L#tT0fL)`-N`N zAZ)QRr513jJ~lPof*W=k;o(P2YQ@wgl;;r;Zag?^s%p4PhNvI@7O|XpZ_1hfdUZCO zx=FFpzYgZ;PJR^ci}!HjnV4__J)OvV3&dO79$yqO+8+&z;s*)J^;qFe;FXq zHSNO@77EGYZzfw;>y!ovSf+yzrED(n_Uwtdo}}pXVn}EkzO=I<&@_LubxvUc{y|f# z1&@lZ7e1enCaTtK5XYMy5X9-qW>$oI%>TlNAW1i)J@|c!`QyR@0T^pH9RW0d1Du^g2 zMc_qUnX(Q~C*jw+ve*hex7?CVRub3sdy{DuAA0 zZ03_nlcoNOsK>0J65WFzFQLn9ysB#QTmpJIXnyCTX29P(jdvlU0UXtdx&R(RUN&AZl5-Uf!j&l49Fuy^N=q;1`@F%1f(YTk@%;Nfi#B9LQ%wT z-~!bS_PS2j@90XaGrd%%IEyA@=-Yb_(ie!E6c0>-yb6+r`+UibRvBR6SXkXWp!vOs@?5j4;q2S@jODpZ&rOy43vKX*APveih1BS%krXmpwP}|3Bc@Aeo3qPit8L)0UvZ8=| zLF|29UUrc6(0`sd9eahsS;}u-K7_hFxJ*S|X295MnSZIy7HaPUceH^y7U6xz!~m!E zswK!O&j0ziI~ql+V8*HgR2-^hr`)T4R*9^0PwZEn9VLIX^6z*shcFrsPE%AjLp;DB zp0zsTp8Vr{;Nwn)Y82fut_^`?N2^s_nq7T;~39Ma0G2UaR-n)47g!HkIV80X;nBzo+Wlmf= zGVFk;&!NBuJx0}}F%lG(h|LaUYRFcBn#zeM6#FGd@f4V`@p{rZk=jTH2jlaEG8Spp z(E*`UfKW&v6z?&iV80YXG5ds2CZcBPsKDvVoi#| z^-n8k{;iv_IBPF*An;8a-i}3f@yDOv!pjif7aN2nhYWu&-v7%3r&HYGsy-`gDm(T? zQ$k9%jkbPt9siVT1^{xRa9a%D##=|j>6GL;YiKQ$%l)LjseZ6rK&;|twq!hdW@4i# zZ^b;j=}uQ~+6+29ycFHkFKI}aI;Un=Ei=V9k+1EQ_wxYltaYlm&dOczg(%At}n06U*`TnI5T+l z&mS7e>n;%{2YMY^_u`Z(Lp$N^`?3=UaShw6Wk%EX{JY^`Zx1~oVoBG`<>I#iEZ%l? zsIE9|%BPSHxmInz#L5q!Hkf+DM=K1=o(Fi;5F-2jLx{zFI+82-&^)uOvD~iZ;J*)$ zPh~P?!q;{y#D}fC)y>TE(Q^?;dCm@|;q;je=ncCT>7^SX>0SfK$v?mUxDhfKEgH+P=2&?t!jc=n1i!IOxoKw570`AixkBHzNI#zrcCCSyDSEs{I)P4 zV0|&%k&J_ox`#0bZ3#CcLg&3Q`;Y)j*Q$dxEkhsCW?k=O?I7Fdw9|$`0U;VsFZr0y zZk8I#dfK!e7MpNTOGD1)IV4CDtadM4tQb1E#W#EC>lI~1qE{X-hn_TNx(9#NuvNzq zGW_0IC7ssx%D)pw6bxTgf0;PGXiWR>qOnL{mB`^aldm4OgMa5X(sO(2`LI--5JP;I ztG_i24C>BOv_k60&sTE?yS)!uek&=DgSdDh#KoVub8`-mxhQJSB7M0BA8u*VfSw~m zvT0VmSb)~X;kv3%s7KUi3QE#bWhLbDxZAaNpMBl44M22PR==#5NvYJ<@#qnJrAznc z;SbgyXKb~1YFX%m#ogA#f!p^B;!MV4hYIkqams^Q_5i<=5=ABPMyChv(oc^f( z&UsaMim_FXXNi$_wzn~}fHwoJp}s{IH#lriIr&uEl6i_+e0ZPwW}n>ei&fii;08hj zD|Jgegh?}9NT{~4`Qfwi_}HW=gGF03Uv0a<`Zv`ba&ptsdc!0Iu`Z-NRMx+xZnZ;a zJ8z;A>!VbC1LDyISUTSph=PU=#Vv7quoUMcVTH>swLR#3SwD0UM@+1F8I~1I->8$m z@Bw1s`ZgqLDBfCZ)0cQbgSfOMdA|`s0lyjWw(*;W_DMoXgA7nZ#a^4+>VJe>FV zu`tDCl~vrkSCyz`W^wVkC0){iavj~0?31a4$iJ*ZTsBz}W8k_0uF-!%?0l0Y_rVnK zPK9ldEfnb0*+{Fjn8ny^lbi&hc+*+4u5axo@yt#J>{)fMxB$;@?;#xzx1=N%PpbZH zR(C;4v1&`CPI^E8azc~XKUcU!^ECK{Ir81^z#Wv`m9l`I1-4c#j%#zi_vJMgjjHuF%9ww(Dr~rnHogf{!N<>A2d^!C> z_nrn-5u<1LgLrG4xQJ2^*xdg{6u#23rsK;(EI=CGoEFkxpMgMXM57BB)%>z1O2JV3JP0tCkSUNuz=i-IN>-oyGMnHKXtb zb5|R>ci#>$rrVeuL=L>_EJ*g}O7k8g?IM14O>n>OF_kGtMo1Fxq+avr-ur+k5dl$Rz$es&l&L0` z&|$yD4iCV>;&o+BLE*BFx<|mR>nC)6k`SLB9iFR08tXW9obn|Q0T$Lj)2%Df-^c!qSvtQ6-kAH8~t+^ zWa=e#8%T_H+(X)nWKe4!!)LIiP-v{l7=Uc`GkoGE6^C|N7UZQ17#vyFbmYuIG6u91 zc@Ad`&bnC-+PnQt!!WZUjZx}>M>O9te^X72l+GL@IHYBZ5Juu?X1e~Kx2rQg6j6-n z30d;zzUO%drqWn`hTE_K%cQMzV3wD1p+u1iaZ0G21sTs6ft7sVJZ-7eN4 zQi}F;CM4o+;V1kh!?>i^eHh2VoqFnY1Q((XfMZ52^C$KA>BuaK`YJrK>g8-; zD}I-nl_f%2!v)2r3vg0pXx1;W<{+LLJ@C*ceUzR%%Eooc4#P)9qV_*_LdN?ANae8A z(4guVe@lT>d!6~h*?Z>74<+!&_#=EM-1Mm}IHSniNG_OJ>D-FFU?fFF=LIkM_bo7g z)j=we8|5Q%k6bs(x41>Rj$~nUyuPt|xti^y|C6}!qH#*Ya}YoF6>h-MWNcL5(u$>R zymFrdgiDk5`T59KM;^3&hP7m>ItBO{2eybg{7ddV+ za)io%%CT3RD++Kag95=u0YrkO5NxE3CAQ0d$#b}w6H^9c2Pbk3{$o?1WMVsaY;lVr zv0WB=AF0Jov=ISbR0Qg=UkZpX2BbCaOc(CN#CFQ(CbkD-rSEdGAz%*(WfPEN#lT4s z0lOOf$M2^#!|?F23Ysn&)JcJ> zTPMl{8pr2Nr9#9q1Fjl}5>GkF+yiEFwyJM3Weau7a;LPrtB_I@PH-lR;vmJ_ zZJCNVHta*os$z719(b0Krk_^OKe*mzQnJn6i{8|RhC-Q<_1%l#wea`H4rzqV5+@qV z>Sl{M`%uTkhyVv1yd`I)FB1Hk`f_l6`SwX-!U`%oCkk>ftcCR%@;4p$m6g)_>~5uM zff3?UzBlVp>2}>}*k{b1EwteyXL{EX_trtT!7`OR+q=#iKdi5e>vX8HZ zm=?M1Uq1dCV#4w$??0^?-ww#K+#|@1W-?(3^OX}0^hIMuKAOa54`mYb1StAj%+=N5 zHTh^UpDo%%tQ%&cY`9~`-S?I`fbS>uD318f`;jIt{D5C8)DOmse=@)mSb@%F>t2|~ zbWvVT>I}psm{Mbn>(GB|yd{5I>!ss=Di{4m?+(Sz%(ZF~;}wU{FZuOMKRRP+f#a(hPWD)z^c`06W0R{b+Zr2@1G-;C>W>+$Z49`wf>TjJ93v5*SRgt~aUP068 z@fWVMy+#$d5=(`K0VmgAfC1nbl6WiLg`UN@28@teL7KvWYzKJ#a;;d&nXIctllBFx z>p#rsCkkB2I_C+mx8qLQ7K;t@FU6-bx*4cAoRCpq43p%{^z0%Rd+tYj8JW+FH-L&4 zJk-8XCbBTe|DwXDoxUQ9Ht5xAGGp5KfuX?us&u@WBK&3}!LLSWBCaG>(hR`&JEsF4Wpkm22ZDePr zrX@G`(wfFOy2sn1&&`TI_LWUaIJI8EwAu{Cok=BzW|Mmg`m&8f-<^2Qs751(&qt;J z?{ziAY)C78Id!IYSu(Qf5VvT+{dsZEdW(j-CRs`OBS9P%ptM=yM*7prQV>(-jXw7_ zq&sxa&V`e3_WP<87MXShS9FV{&4!O1h;#h*REPzR#g%4kMn6qpIsL<;XO|gUKoZt) zv@%GdL+i~v4UG9Jxy0=K!*`$?G9aGs#{afo|Cg07o9@?tlp0ylwhQM3Bx!qJd+$1(zcgd9tvuP^n<=0P#CrGj z4u`+VM?)uf1y^_PM~70bszm7<(0aMwxMe9;RW2))^d+ag%FJRg8OAW=+7mYed(X1? zq7;)wwp))i&dzsIH*@KGL_WMpqY3yd%f+^rJsZQY+WvjzjW6c;b4(ohc{Ia<4C!Ig zYiAXo4O*7OoFbDhvD=G8OcSR2ImO22^oky6uzEc6LZ=})3a@VNaD#@nTkJhOv~L|H z3B&!Bj`WH~i%fpOYV~6}v93KI(*2+oA>N&L9+%&g!iIB-jNar4r{gyBNG08D*_H4k zu?^%{+=B|Kx3{3s*JZz@eZ@60zr(P>zgd5Mj+tYa{wt{NHp@+B`bTN&&Is%z($C4AlW&NtP_ zwQ!$~*z7=2Gd6UhA>Bz(sqbCrTZ+qL< ztC;(xxk+(6vHI!h!N=Mn#@#cjEw{8qp?1YJ;*B`IqzI}e>|Aq2M*9XY@<@aByy3Iy z{yIR9cy$l%-^vEpRB_O9(Xb@Viu!Z_P`@M#kvK0R2z6mqn0 z2tTT+F_zu-4Lvhy`#BBtjv@n^|>WAZh7H!O3+Lodkpx}GwI%$NP81M$6F}G>gOI{^&JIC z;rNXm2Qn&8=SILN{k(b}4VEKxW+1hAo1re+k?#@DB4a;-Pf|TNCB>9gw7tBWqF!LB z^vnXDUsCOoz44rCn+JzXV>qJ9av;xT*|&yf?P9F{AI>xF^x#TrezL}@MVJGr=uZ#7 z3`a}(DMayGO1R&@i_v$S=Jb$G=9=!((>8+9V*qLx{luhie@FJj;^%R{{8sHr+oDol z$nY?)y{WNt%=3~mhXgbE&1J7zAW!5V5i+k;yBXq6AbQ+IZ_<<1Y%j7)h`^I12&@al zBMC8s{g07fABwT_uC@r+LbfAgw3N^*|ND@Q`(hLF=id%Tv+6RC_91@kbVNb)l>ANf zC^BN2**TCAlTQczzJO)YL&6HBq$pzVj+yd4b%;qBzA-`JfgI!Cibu~vWJT$(loMpz(TPx2B%Pp1xasVmG%BpL%+S*+QVFV zO)E@WL~|A8X+aGP4LGcS)WE2U(zB5H?U~#SDh@AXZ~`LwH-sU_XcX<$l||L{M*zyq zi6&2pJB6}QIOvG$G=PyijwsL^i~>nfD#I zW(weJ@wJdCaSFDFo8DUN94^Mo;NXLXbSjmkp*X)G(B7-nFRS0VT zTnlp6y|a)Pi>Ytnr*S(Fo&jj3R+x?n)ZiSopmu3Ok}h^`BAgiNYsUGl$$l-yM%~F$ zb-pZ(eUS*LT9GqUKKp{IbEWouK5dsum5=(~=}2|8gF05c;|_=5lHjPnks=h&A^l%s ztuPgxg-v9u|Ce41sce79jLK;hXEi=C*;DX6Oym-jN}7ecC;*I{Ch}S?r?YsE$;G(E z9LSck;B#!{bvjpt98>oUAFO*hH0`~t4TO?DkXd_Gs!r(CO+{w8{9Q-nK2}F$+J@B- ztyRb79*(T#KeuG0)~Bz{2^=KGF4heDA$+Cri!|iua_Vudupg*TeCPjx9Bt2pgrkhD z>O!pkXc^D_cRe**`B54CvD7HpDKq@W=mzmOVN2a=;^F(4e3{zARDC}#oOq=OtNR1~ zR$$D^i5L?O%^IJ7fz@q(7REd{X6cxkyO9P{sOLxlBI4i}5#cyWM0_H1i39mE$d)ww#NCF{|9&DU3mZi literal 0 HcmV?d00001 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