Implement candidate intersections

CHANGELOG.md

@@ -101,7 +101,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148]
 - Support 64-bit hex literals and unary operations in constants [#6616](https://github.com/gfx-rs/wgpu/pull/6616).
 - Implement `quantizeToF16()` for WGSL frontend, and WGSL, SPIR-V, HLSL, MSL, and GLSL backends. By @jamienicol in [#6519](https://github.com/gfx-rs/wgpu/pull/6519).
 - Add support for GLSL `usampler*` and `isampler*`. By @DavidPeicho in [#6513](https://github.com/gfx-rs/wgpu/pull/6513).
-- Expose Ray Query flags as constants in WGSL. By @kvark in [#5429](https://github.com/gfx-rs/wgpu/pull/5429)
+- Expose Ray Query flags as constants in WGSL. Implement candidate intersections. By @kvark in [#5429](https://github.com/gfx-rs/wgpu/pull/5429)
 
 #### General
 

examples/src/ray_cube_compute/shader.wgsl

@@ -15,7 +15,7 @@ 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;
+let RAY_QUERY_INTERSECTION_AABB = 3u;
 
 struct RayDesc {
     flags: u32,
@@ -60,7 +60,6 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     let target_size = textureDimensions(output);
     var color =  vec4<f32>(vec2<f32>(global_id.xy) / vec2<f32>(target_size), 0.0, 1.0);
 
-
 	let pixel_center = vec2<f32>(global_id.xy) + vec2<f32>(0.5);
 	let in_uv = pixel_center/vec2<f32>(target_size.xy);
 	let d = in_uv * 2.0 - 1.0;

examples/src/ray_scene/shader.wgsl

@@ -38,7 +38,7 @@ 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;
+let RAY_QUERY_INTERSECTION_AABB = 3u;
 
 struct RayDesc {
     flags: u32,

naga/src/back/msl/writer.rs

@@ -2254,14 +2254,14 @@ impl<W: Write> Writer<W> {
                     write!(self.out, ")")?;
                 }
             }
-            crate::Expression::RayQueryGetIntersection { query, committed } => {
+            crate::Expression::RayQueryGetIntersection {
+                query,
+                committed: _,
+            } => {
                 if context.lang_version < (2, 4) {
                     return Err(Error::UnsupportedRayTracing);
                 }
 
-                if !committed {
-                    unimplemented!()
-                }
                 let ty = context.module.special_types.ray_intersection.unwrap();
                 let type_name = &self.names[&NameKey::Type(ty)];
                 write!(self.out, "{type_name} {{{RAY_QUERY_FUN_MAP_INTERSECTION}(")?;

naga/src/back/spv/block.rs

@@ -1736,10 +1736,7 @@ impl BlockContext<'_> {
             }
             crate::Expression::ArrayLength(expr) => self.write_runtime_array_length(expr, block)?,
             crate::Expression::RayQueryGetIntersection { query, committed } => {
-                if !committed {
-                    return Err(Error::FeatureNotImplemented("candidate intersection"));
-                }
-                self.write_ray_query_get_intersection(query, block)
+                self.write_ray_query_get_intersection(query, block, committed)
             }
         };
 

naga/src/back/spv/ray.rs

@@ -106,23 +106,60 @@ impl BlockContext<'_> {
         &mut self,
         query: Handle<crate::Expression>,
         block: &mut Block,
+        is_committed: bool,
     ) -> spirv::Word {
         let query_id = self.cached[query];
-        let intersection_id = self.writer.get_constant_scalar(crate::Literal::U32(
-            spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR as _,
-        ));
+        let intersection_id =
+            self.writer
+                .get_constant_scalar(crate::Literal::U32(if is_committed {
+                    spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR
+                } else {
+                    spirv::RayQueryIntersection::RayQueryCandidateIntersectionKHR
+                } as _));
 
         let flag_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
             NumericType::Scalar(crate::Scalar::U32),
         )));
-        let kind_id = self.gen_id();
+        let raw_kind_id = self.gen_id();
         block.body.push(Instruction::ray_query_get_intersection(
             spirv::Op::RayQueryGetIntersectionTypeKHR,
             flag_type_id,
-            kind_id,
+            raw_kind_id,
             query_id,
             intersection_id,
         ));
+        let kind_id = if is_committed {
+            // Nothing to do: the IR value matches `spirv::RayQueryCommittedIntersectionType`
+            raw_kind_id
+        } else {
+            // Remap from the candidate kind to IR
+            let condition_id = self.gen_id();
+            let committed_triangle_kind_id = self.writer.get_constant_scalar(crate::Literal::U32(
+                spirv::RayQueryCandidateIntersectionType::RayQueryCandidateIntersectionTriangleKHR
+                    as _,
+            ));
+            block.body.push(Instruction::binary(
+                spirv::Op::IEqual,
+                self.writer.get_bool_type_id(),
+                condition_id,
+                raw_kind_id,
+                committed_triangle_kind_id,
+            ));
+            let kind_id = self.gen_id();
+            block.body.push(Instruction::select(
+                flag_type_id,
+                kind_id,
+                condition_id,
+                self.writer.get_constant_scalar(crate::Literal::U32(
+                    crate::RayQueryIntersection::Triangle as _,
+                )),
+                self.writer.get_constant_scalar(crate::Literal::U32(
+                    crate::RayQueryIntersection::Aabb as _,
+                )),
+            ));
+            kind_id
+        };
+
         let instance_custom_index_id = self.gen_id();
         block.body.push(Instruction::ray_query_get_intersection(
             spirv::Op::RayQueryGetIntersectionInstanceCustomIndexKHR,
@@ -201,6 +238,8 @@ impl BlockContext<'_> {
             query_id,
             intersection_id,
         ));
+        //Note: there is also `OpRayQueryGetIntersectionCandidateAABBOpaqueKHR`,
+        // but it's not a property of an intersection.
 
         let transform_type_id =
             self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Matrix {

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

@@ -2568,12 +2568,22 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                             args.finish()?;
 
                             let _ = ctx.module.generate_ray_intersection_type();
-
                             crate::Expression::RayQueryGetIntersection {
                                 query,
                                 committed: true,
                             }
                         }
+                        "rayQueryGetCandidateIntersection" => {
+                            let mut args = ctx.prepare_args(arguments, 1, span);
+                            let query = self.ray_query_pointer(args.next()?, ctx)?;
+                            args.finish()?;
+
+                            let _ = ctx.module.generate_ray_intersection_type();
+                            crate::Expression::RayQueryGetIntersection {
+                                query,
+                                committed: false,
+                            }
+                        }
                         "RayDesc" => {
                             let ty = ctx.module.generate_ray_desc_type();
                             let handle = self.construct(

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

@@ -661,7 +661,7 @@ impl Parser {
         const fn literal_ray_intersection<'b>(
             intersection: crate::RayQueryIntersection,
         ) -> ast::Expression<'b> {
-            ast::Expression::Literal(ast::Literal::Number(Number::U32(intersection.bits())))
+            ast::Expression::Literal(ast::Literal::Number(Number::U32(intersection as u32)))
         }
 
         let expr = match lexer.peek() {
@@ -739,19 +739,19 @@ impl Parser {
             }
             (Token::Word("RAY_QUERY_INTERSECTION_NONE"), _) => {
                 let _ = lexer.next();
-                literal_ray_intersection(crate::RayQueryIntersection::empty())
+                literal_ray_intersection(crate::RayQueryIntersection::None)
             }
             (Token::Word("RAY_QUERY_INTERSECTION_TRIANGLE"), _) => {
                 let _ = lexer.next();
-                literal_ray_intersection(crate::RayQueryIntersection::TRIANGLE)
+                literal_ray_intersection(crate::RayQueryIntersection::Triangle)
             }
             (Token::Word("RAY_QUERY_INTERSECTION_GENERATED"), _) => {
                 let _ = lexer.next();
-                literal_ray_intersection(crate::RayQueryIntersection::GENERATED)
+                literal_ray_intersection(crate::RayQueryIntersection::Generated)
             }
             (Token::Word("RAY_QUERY_INTERSECTION_AABB"), _) => {
                 let _ = lexer.next();
-                literal_ray_intersection(crate::RayQueryIntersection::AABB)
+                literal_ray_intersection(crate::RayQueryIntersection::Aabb)
             }
             (Token::Word(word), span) => {
                 let start = lexer.start_byte_offset();

naga/src/lib.rs

@@ -2328,23 +2328,26 @@ bitflags::bitflags! {
     }
 }
 
-bitflags::bitflags! {
-    /// Type of a ray query intersection.
-    /// Matching vulkan constants can be found in
-    /// https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/KHR/SPV_KHR_ray_query.asciidoc
-    #[cfg_attr(feature = "serialize", derive(Serialize))]
-    #[cfg_attr(feature = "deserialize", derive(Deserialize))]
-    #[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
-    #[derive(Clone, Copy, Debug, Default, Eq, Hash, Ord, PartialEq, PartialOrd)]
-    pub struct RayQueryIntersection: u32 {
-        /// Intersecting with triangles..
-        /// Matches RayQueryCommittedIntersectionTriangleKHR and RayQueryCandidateIntersectionTriangleKHR.
-        const TRIANGLE = 0x1;
-        /// Intersecting with generated primitives.
-        /// Matches RayQueryCommittedIntersectionGeneratedKHR.
-        const GENERATED = 0x2;
-        /// Intersecting with Axis Aligned Bounding Boxes.
-        /// Matches RayQueryCandidateIntersectionAABBKHR.
-        const AABB = 0x4;
-    }
+/// Type of a ray query intersection.
+/// Matching vulkan constants can be found in
+/// <https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/KHR/SPV_KHR_ray_query.asciidoc>
+/// but the actual values are different for candidate intersections.
+#[cfg_attr(feature = "serialize", derive(Serialize))]
+#[cfg_attr(feature = "deserialize", derive(Deserialize))]
+#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
+#[derive(Clone, Copy, Debug, Default, Eq, Hash, Ord, PartialEq, PartialOrd)]
+pub enum RayQueryIntersection {
+    /// No intersection found.
+    /// Matches `RayQueryCommittedIntersectionNoneKHR`.
+    #[default]
+    None = 0,
+    /// Intersecting with triangles.
+    /// Matches `RayQueryCommittedIntersectionTriangleKHR` and `RayQueryCandidateIntersectionTriangleKHR`.
+    Triangle = 1,
+    /// Intersecting with generated primitives.
+    /// Matches `RayQueryCommittedIntersectionGeneratedKHR`.
+    Generated = 2,
+    /// Intersecting with Axis Aligned Bounding Boxes.
+    /// Matches `RayQueryCandidateIntersectionAABBKHR`.
+    Aabb = 3,
 }

naga/tests/in/ray-query.wgsl

@@ -14,7 +14,7 @@ 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;
+let RAY_QUERY_INTERSECTION_AABB = 3u;
 
 struct RayDesc {
     flags: u32,
@@ -78,3 +78,14 @@ fn main() {
     output.visible = u32(intersection.kind == RAY_QUERY_INTERSECTION_NONE);
     output.normal = get_torus_normal(dir * intersection.t, intersection);
 }
+
+@compute @workgroup_size(1)
+fn main_candidate() {
+    let pos = vec3<f32>(0.0);
+    let dir = vec3<f32>(0.0, 1.0, 0.0);
+
+    var rq: ray_query;
+    rayQueryInitialize(&rq, acc_struct, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, 0xFFu, 0.1, 100.0, pos, dir));
+    let intersection = rayQueryGetCandidateIntersection(&rq);
+    output.visible = u32(intersection.kind == RAY_QUERY_INTERSECTION_AABB);
+}

naga/tests/out/msl/ray-query.msl

@@ -46,24 +46,24 @@ RayIntersection query_loop(
     metal::float3 dir,
     metal::raytracing::instance_acceleration_structure acs
 ) {
-    _RayQuery rq = {};
+    _RayQuery rq_1 = {};
     RayDesc _e8 = RayDesc {4u, 255u, 0.1, 100.0, pos, dir};
-    rq.intersector.assume_geometry_type(metal::raytracing::geometry_type::triangle);
-    rq.intersector.set_opacity_cull_mode((_e8.flags & 64) != 0 ? metal::raytracing::opacity_cull_mode::opaque : (_e8.flags & 128) != 0 ? metal::raytracing::opacity_cull_mode::non_opaque : metal::raytracing::opacity_cull_mode::none);
-    rq.intersector.force_opacity((_e8.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e8.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none);
-    rq.intersector.accept_any_intersection((_e8.flags & 4) != 0);
-    rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e8.origin, _e8.dir, _e8.tmin, _e8.tmax), acs, _e8.cull_mask);    rq.ready = true;
+    rq_1.intersector.assume_geometry_type(metal::raytracing::geometry_type::triangle);
+    rq_1.intersector.set_opacity_cull_mode((_e8.flags & 64) != 0 ? metal::raytracing::opacity_cull_mode::opaque : (_e8.flags & 128) != 0 ? metal::raytracing::opacity_cull_mode::non_opaque : metal::raytracing::opacity_cull_mode::none);
+    rq_1.intersector.force_opacity((_e8.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e8.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none);
+    rq_1.intersector.accept_any_intersection((_e8.flags & 4) != 0);
+    rq_1.intersection = rq_1.intersector.intersect(metal::raytracing::ray(_e8.origin, _e8.dir, _e8.tmin, _e8.tmax), acs, _e8.cull_mask);    rq_1.ready = true;
     while(true) {
-        bool _e9 = rq.ready;
-        rq.ready = false;
+        bool _e9 = rq_1.ready;
+        rq_1.ready = false;
         if (_e9) {
         } else {
             break;
         }
 #define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
         LOOP_IS_BOUNDED
     }
-    return RayIntersection {_map_intersection_type(rq.intersection.type), rq.intersection.distance, rq.intersection.user_instance_id, rq.intersection.instance_id, {}, rq.intersection.geometry_id, rq.intersection.primitive_id, rq.intersection.triangle_barycentric_coord, rq.intersection.triangle_front_facing, {}, rq.intersection.object_to_world_transform, rq.intersection.world_to_object_transform};
+    return RayIntersection {_map_intersection_type(rq_1.intersection.type), rq_1.intersection.distance, rq_1.intersection.user_instance_id, rq_1.intersection.instance_id, {}, rq_1.intersection.geometry_id, rq_1.intersection.primitive_id, rq_1.intersection.triangle_barycentric_coord, rq_1.intersection.triangle_front_facing, {}, rq_1.intersection.object_to_world_transform, rq_1.intersection.world_to_object_transform};
 }
 
 metal::float3 get_torus_normal(
@@ -88,3 +88,22 @@ kernel void main_(
     output.normal = _e18;
     return;
 }
+
+
+kernel void main_candidate(
+  metal::raytracing::instance_acceleration_structure acc_struct [[user(fake0)]]
+, device Output& output [[user(fake0)]]
+) {
+    _RayQuery rq = {};
+    metal::float3 pos_2 = metal::float3(0.0);
+    metal::float3 dir_2 = metal::float3(0.0, 1.0, 0.0);
+    RayDesc _e12 = RayDesc {4u, 255u, 0.1, 100.0, pos_2, dir_2};
+    rq.intersector.assume_geometry_type(metal::raytracing::geometry_type::triangle);
+    rq.intersector.set_opacity_cull_mode((_e12.flags & 64) != 0 ? metal::raytracing::opacity_cull_mode::opaque : (_e12.flags & 128) != 0 ? metal::raytracing::opacity_cull_mode::non_opaque : metal::raytracing::opacity_cull_mode::none);
+    rq.intersector.force_opacity((_e12.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e12.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none);
+    rq.intersector.accept_any_intersection((_e12.flags & 4) != 0);
+    rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e12.origin, _e12.dir, _e12.tmin, _e12.tmax), acc_struct, _e12.cull_mask);    rq.ready = true;
+    RayIntersection intersection_1 = RayIntersection {_map_intersection_type(rq.intersection.type), rq.intersection.distance, rq.intersection.user_instance_id, rq.intersection.instance_id, {}, rq.intersection.geometry_id, rq.intersection.primitive_id, rq.intersection.triangle_barycentric_coord, rq.intersection.triangle_front_facing, {}, rq.intersection.object_to_world_transform, rq.intersection.world_to_object_transform};
+    output.visible = static_cast<uint>(intersection_1.kind == 3u);
+    return;
+}

naga/tests/out/spv/ray-query.spvasm

@@ -1,14 +1,16 @@
 ; SPIR-V
 ; Version: 1.4
 ; Generator: rspirv
-; Bound: 104
+; Bound: 136
 OpCapability Shader
 OpCapability RayQueryKHR
 OpExtension "SPV_KHR_ray_query"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
 OpEntryPoint GLCompute %84 "main" %15 %17
+OpEntryPoint GLCompute %105 "main_candidate" %15 %17
 OpExecutionMode %84 LocalSize 1 1 1
+OpExecutionMode %105 LocalSize 1 1 1
 OpMemberDecorate %10 0 Offset 0
 OpMemberDecorate %10 1 Offset 4
 OpMemberDecorate %10 2 Offset 8
@@ -74,6 +76,8 @@ OpMemberDecorate %18 0 Offset 0
 %91 = OpConstantComposite  %4  %70 %68 %70
 %94 = OpTypePointer StorageBuffer %6
 %99 = OpTypePointer StorageBuffer %4
+%108 = OpConstantComposite  %12  %27 %28 %29 %30 %90 %91
+%109 = OpConstant  %6  3
 %25 = OpFunction  %10  None %26
 %21 = OpFunctionParameter  %4
 %22 = OpFunctionParameter  %4
@@ -161,4 +165,39 @@ OpStore %98 %97
 %103 = OpAccessChain  %99  %89 %50
 OpStore %103 %102
 OpReturn
+OpFunctionEnd
+%105 = OpFunction  %2  None %85
+%104 = OpLabel
+%110 = OpVariable  %32  Function
+%106 = OpLoad  %5  %15
+%107 = OpAccessChain  %87  %17 %88
+OpBranch %111
+%111 = OpLabel
+%112 = OpCompositeExtract  %6  %108 0
+%113 = OpCompositeExtract  %6  %108 1
+%114 = OpCompositeExtract  %3  %108 2
+%115 = OpCompositeExtract  %3  %108 3
+%116 = OpCompositeExtract  %4  %108 4
+%117 = OpCompositeExtract  %4  %108 5
+OpRayQueryInitializeKHR %110 %106 %112 %113 %116 %114 %117 %115
+%118 = OpRayQueryGetIntersectionTypeKHR  %6  %110 %88
+%119 = OpIEqual  %8  %118 %88
+%120 = OpSelect  %6  %119 %50 %109
+%121 = OpRayQueryGetIntersectionInstanceCustomIndexKHR  %6  %110 %88
+%122 = OpRayQueryGetIntersectionInstanceIdKHR  %6  %110 %88
+%123 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR  %6  %110 %88
+%124 = OpRayQueryGetIntersectionGeometryIndexKHR  %6  %110 %88
+%125 = OpRayQueryGetIntersectionPrimitiveIndexKHR  %6  %110 %88
+%126 = OpRayQueryGetIntersectionTKHR  %3  %110 %88
+%127 = OpRayQueryGetIntersectionBarycentricsKHR  %7  %110 %88
+%128 = OpRayQueryGetIntersectionFrontFaceKHR  %8  %110 %88
+%129 = OpRayQueryGetIntersectionObjectToWorldKHR  %9  %110 %88
+%130 = OpRayQueryGetIntersectionWorldToObjectKHR  %9  %110 %88
+%131 = OpCompositeConstruct  %10  %120 %126 %121 %122 %123 %124 %125 %127 %128 %129 %130
+%132 = OpCompositeExtract  %6  %131 0
+%133 = OpIEqual  %8  %132 %109
+%134 = OpSelect  %6  %133 %50 %88
+%135 = OpAccessChain  %94  %107 %88
+OpStore %135 %134
+OpReturn
 OpFunctionEnd