Repository: gfx-rs/wgpu (language: rust)

Issue:
Allow for comitting candidate intersections
**Is your feature request related to a problem? Please describe.**
Currently Ray Queries support getting candidate intersections. However unlike spirv, naga does not have a way to commit these candidate intersection.

**Describe the solution you'd like**
A way to commit the candidate intersections like `OpRayQueryConfirmIntersectionKHR`

**Describe alternatives you've considered**
Just let the user use and compare it with the committed intersections - this is probably slower and uses more memory.

**Additional context**
N/A



Discussion / hints:
Awesome! I am looking forward for this implementation. I am 100% confident this is necessary for user custom intersections (aka procedural or AABB) in the future.

Fix patch:
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 726efb6ade..839af9d6a3 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -45,6 +45,7 @@ Bottom level categories:
 #### Naga
 
 - Support @must_use attribute on function declarations. By @turbocrime in [#6801](https://github.com/gfx-rs/wgpu/pull/6801).
+- Support for generating the candidate intersections from AABB geometry, and confirming the hits. By @kvark in [#7047](https://github.com/gfx-rs/wgpu/pull/7047).
 
 ### Changes
 
diff --git a/etc/specs/ray_tracing.md b/etc/specs/ray_tracing.md
index f23b5305a2..76ce6ce15a 100644
--- a/etc/specs/ray_tracing.md
+++ b/etc/specs/ray_tracing.md
@@ -2,7 +2,7 @@
 
 🧪Experimental🧪
 
-`wgpu` supports an experimental version of ray tracing which is subject to change. The extensions allow for acceleration structures to be created and built (with 
+`wgpu` supports an experimental version of ray tracing which is subject to change. The extensions allow for acceleration structures to be created and built (with
 `Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE` enabled) and interacted with in shaders. Currently `naga` only supports ray queries
 (accessible with `Features::EXPERIMENTAL_RAY_QUERY` enabled in wgpu).
 
@@ -10,15 +10,15 @@
 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/1040).
 Large changes may mean that this documentation may be out of date.
 
-***This is not*** an introduction to raytracing, and assumes basic prior knowledge, to look at the fundamentals look at 
+***This is not*** an introduction to raytracing, and assumes basic prior knowledge, to look at the fundamentals look at
 an [introduction](https://developer.nvidia.com/blog/introduction-nvidia-rtx-directx-ray-tracing/).
 
 ## `wgpu`'s raytracing API:
 
 The documentation and specific details of the functions and structures provided
-can be found with their definitions.  
+can be found with their definitions.
 
-A [`Blas`] can be created with [`Device::create_blas`].  
+A [`Blas`] can be created with [`Device::create_blas`].
 A [`Tlas`] can be created with [`Device::create_tlas`].
 
 Unless one is planning on using the unsafe building API (not recommended for beginners) a [`Tlas`] should be put inside
@@ -57,12 +57,21 @@ rayQueryInitialize(rq: ptr<function, ray_query>, acceleration_structure: acceler
 // - The hit is considered `Candidate` if this function returns true, and the hit is considered `Committed` if
 //   this function returns false.
 // - A `Candidate` intersection interrupts the ray traversal.
-// - A `Candidate` intersection may happen anywhere along the ray, it should not be relied on to give the closest hit. A 
+// - A `Candidate` intersection may happen anywhere along the ray, it should not be relied on to give the closest hit. A
 //   `Candidate` intersection is to allow the user themselves to decide if that intersection is valid*. If one wants to get
 //   the closest hit a `Committed` intersection should be used.
 // - Calling this function multiple times will cause the ray traversal to continue if it was interrupted by a `Candidate`
 //   intersection.
-rayQueryProceed(rq: ptr<function, ray_query>) -> bool`
+rayQueryProceed(rq: ptr<function, ray_query>) -> bool
+
+// - Generates a hit from procedural geometry at a particular distance.
+rayQueryGenerateIntersection(hit_t: f32)
+
+// - Commits a hit from triangular non-opaque geometry.
+rayQueryConfirmIntersection()
+
+// - Aborts the query.
+rayQueryTerminate()
 
 // - Returns intersection details about a hit considered `Committed`.
 rayQueryGetCommittedIntersection(rq: ptr<function, ray_query>) -> RayIntersection
@@ -71,19 +80,24 @@ rayQueryGetCommittedIntersection(rq: ptr<function, ray_query>) -> RayIntersectio
 rayQueryGetCandidateIntersection(rq: ptr<function, ray_query>) -> RayIntersection
 ```
 
-*The API to commit a candidate intersection is not yet implemented but would be possible to be user implemented.
-
 > [!CAUTION]
-> 
+>
 > ### ⚠️Undefined behavior ⚠️:
 > - Calling `rayQueryGetCommittedIntersection` or `rayQueryGetCandidateIntersection` when `rayQueryProceed` has not been
 > called on this ray query since it was initialized (or if the ray query has not been previously initialized).
 > - Calling `rayQueryGetCommittedIntersection` when `rayQueryProceed`'s latest return on this ray query is considered
->   `Candidate`. 
+>   `Candidate`.
 > - Calling `rayQueryGetCandidateIntersection` when `rayQueryProceed`'s latest return on this ray query is considered
 >   `Committed`.
 > - 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`,
+> - Calling `rayQueryGenerateIntersection` with `hit_t` outside of `RayDesc::t_min .. RayDesc::t_max` range.
+>   or when `rayQueryProceed`'s latest return on this ray query is not considered `Candidate`.
+> - Calling `rayQueryConfirmIntersection` on a query with last intersection kind not being
+>   `RAY_QUERY_INTERSECTION_TRIANGLE`,
+>   or when `rayQueryProceed`'s latest return on this ray query is not considered `Candidate`.
+>
 > *this is only known undefined behaviour, and will be worked around in the future.
 
 ```wgsl
diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs
index e44e8d8eae..67b1b9dc92 100644
--- a/naga/src/back/dot/mod.rs
+++ b/naga/src/back/dot/mod.rs
@@ -293,6 +293,13 @@ impl StatementGraph {
                             self.emits.push((id, result));
                             "RayQueryProceed"
                         }
+                        crate::RayQueryFunction::GenerateIntersection { hit_t } => {
+                            self.dependencies.push((id, hit_t, "hit_t"));
+                            "RayQueryGenerateIntersection"
+                        }
+                        crate::RayQueryFunction::ConfirmIntersection => {
+                            "RayQueryConfirmIntersection"
+                        }
                         crate::RayQueryFunction::Terminate => "RayQueryTerminate",
                     }
                 }
diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs
index 7b6826579b..ee0473573e 100644
--- a/naga/src/back/hlsl/writer.rs
+++ b/naga/src/back/hlsl/writer.rs
@@ -2452,7 +2452,20 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     self.write_expr(module, query, func_ctx)?;
                     writeln!(self.out, ".Proceed();")?;
                 }
+                RayQueryFunction::GenerateIntersection { hit_t } => {
+                    write!(self.out, "{level}")?;
+                    self.write_expr(module, query, func_ctx)?;
+                    write!(self.out, ".CommitProceduralPrimitiveHit(")?;
+                    self.write_expr(module, hit_t, func_ctx)?;
+                    writeln!(self.out, ");")?;
+                }
+                RayQueryFunction::ConfirmIntersection => {
+                    write!(self.out, "{level}")?;
+                    self.write_expr(module, query, func_ctx)?;
+                    writeln!(self.out, ".CommitNonOpaqueTriangleHit();")?;
+                }
                 RayQueryFunction::Terminate => {
+                    write!(self.out, "{level}")?;
                     self.write_expr(module, query, func_ctx)?;
                     writeln!(self.out, ".Abort();")?;
                 }
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 402cb292ef..cc58693f55 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -30,6 +30,7 @@ const RT_NAMESPACE: &str = "metal::raytracing";
 const RAY_QUERY_TYPE: &str = "_RayQuery";
 const RAY_QUERY_FIELD_INTERSECTOR: &str = "intersector";
 const RAY_QUERY_FIELD_INTERSECTION: &str = "intersection";
+const RAY_QUERY_MODERN_SUPPORT: bool = false; //TODO
 const RAY_QUERY_FIELD_READY: &str = "ready";
 const RAY_QUERY_FUN_MAP_INTERSECTION: &str = "_map_intersection_type";
 
@@ -3506,16 +3507,41 @@ impl<W: Write> Writer<W> {
                             self.named_expressions.insert(result, name);
                             self.put_expression(query, &context.expression, true)?;
                             writeln!(self.out, ".{RAY_QUERY_FIELD_READY};")?;
-                            //TODO: actually proceed?
-
-                            write!(self.out, "{level}")?;
-                            self.put_expression(query, &context.expression, true)?;
-                            writeln!(self.out, ".{RAY_QUERY_FIELD_READY} = false;")?;
+                            if RAY_QUERY_MODERN_SUPPORT {
+                                write!(self.out, "{level}")?;
+                                self.put_expression(query, &context.expression, true)?;
+                                writeln!(self.out, ".?.next();")?;
+                            }
+                        }
+                        crate::RayQueryFunction::GenerateIntersection { hit_t } => {
+                            if RAY_QUERY_MODERN_SUPPORT {
+                                write!(self.out, "{level}")?;
+                                self.put_expression(query, &context.expression, true)?;
+                                write!(self.out, ".?.commit_bounding_box_intersection(")?;
+                                self.put_expression(hit_t, &context.expression, true)?;
+                                writeln!(self.out, ");")?;
+                            } else {
+                                log::warn!("Ray Query GenerateIntersection is not yet supported");
+                            }
+                        }
+                        crate::RayQueryFunction::ConfirmIntersection => {
+                            if RAY_QUERY_MODERN_SUPPORT {
+                                write!(self.out, "{level}")?;
+                                self.put_expression(query, &context.expression, true)?;
+                                writeln!(self.out, ".?.commit_triangle_intersection();")?;
+                            } else {
+                                log::warn!("Ray Query ConfirmIntersection is not yet supported");
+                            }
                         }
                         crate::RayQueryFunction::Terminate => {
+                            if RAY_QUERY_MODERN_SUPPORT {
+                                write!(self.out, "{level}")?;
+                                self.put_expression(query, &context.expression, true)?;
+                                writeln!(self.out, ".?.abort();")?;
+                            }
                             write!(self.out, "{level}")?;
                             self.put_expression(query, &context.expression, true)?;
-                            writeln!(self.out, ".{RAY_QUERY_FIELD_INTERSECTION}.abort();")?;
+                            writeln!(self.out, ".{RAY_QUERY_FIELD_READY} = false;")?;
                         }
                     }
                 }
diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs
index bb9fb7f448..05e6bdb374 100644
--- a/naga/src/back/pipeline_constants.rs
+++ b/naga/src/back/pipeline_constants.rs
@@ -821,6 +821,10 @@ fn adjust_stmt(new_pos: &HandleVec<Expression, Handle<Expression>>, stmt: &mut S
                 crate::RayQueryFunction::Proceed { ref mut result } => {
                     adjust(result);
                 }
+                crate::RayQueryFunction::GenerateIntersection { ref mut hit_t } => {
+                    adjust(hit_t);
+                }
+                crate::RayQueryFunction::ConfirmIntersection => {}
                 crate::RayQueryFunction::Terminate => {}
             }
         }
diff --git a/naga/src/back/spv/instructions.rs b/naga/src/back/spv/instructions.rs
index 38aed8c351..57bfb4e9f3 100644
--- a/naga/src/back/spv/instructions.rs
+++ b/naga/src/back/spv/instructions.rs
@@ -779,6 +779,19 @@ impl super::Instruction {
         instruction
     }
 
+    pub(super) fn ray_query_generate_intersection(query: Word, hit: Word) -> Self {
+        let mut instruction = Self::new(Op::RayQueryGenerateIntersectionKHR);
+        instruction.add_operand(query);
+        instruction.add_operand(hit);
+        instruction
+    }
+
+    pub(super) fn ray_query_confirm_intersection(query: Word) -> Self {
+        let mut instruction = Self::new(Op::RayQueryConfirmIntersectionKHR);
+        instruction.add_operand(query);
+        instruction
+    }
+
     pub(super) fn ray_query_get_intersection(
         op: Op,
         result_type_id: Word,
diff --git a/naga/src/back/spv/ray.rs b/naga/src/back/spv/ray.rs
index 907072a267..b8e5281ac9 100644
--- a/naga/src/back/spv/ray.rs
+++ b/naga/src/back/spv/ray.rs
@@ -609,6 +609,19 @@ impl BlockContext<'_> {
                     .body
                     .push(Instruction::ray_query_proceed(result_type_id, id, query_id));
             }
+            crate::RayQueryFunction::GenerateIntersection { hit_t } => {
+                let hit_id = self.cached[hit_t];
+                block
+                    .body
+                    .push(Instruction::ray_query_generate_intersection(
+                        query_id, hit_id,
+                    ));
+            }
+            crate::RayQueryFunction::ConfirmIntersection => {
+                block
+                    .body
+                    .push(Instruction::ray_query_confirm_intersection(query_id));
+            }
             crate::RayQueryFunction::Terminate => {}
         }
     }
diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs
index 596f9d4067..08b1ea9757 100644
--- a/naga/src/compact/statements.rs
+++ b/naga/src/compact/statements.rs
@@ -190,6 +190,10 @@ impl FunctionTracer<'_> {
             Qf::Proceed { result } => {
                 self.expressions_used.insert(result);
             }
+            Qf::GenerateIntersection { hit_t } => {
+                self.expressions_used.insert(hit_t);
+            }
+            Qf::ConfirmIntersection => {}
             Qf::Terminate => {}
         }
     }
@@ -393,6 +397,10 @@ impl FunctionMap {
             Qf::Proceed { ref mut result } => {
                 self.expressions.adjust(result);
             }
+            Qf::GenerateIntersection { ref mut hit_t } => {
+                self.expressions.adjust(hit_t);
+            }
+            Qf::ConfirmIntersection => {}
             Qf::Terminate => {}
         }
     }
diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs
index d25eb362c1..833b24c270 100644
--- a/naga/src/front/wgsl/lower/mod.rs
+++ b/naga/src/front/wgsl/lower/mod.rs
@@ -2692,6 +2692,40 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                                 .push(crate::Statement::RayQuery { query, fun }, span);
                             return Ok(Some(result));
                         }
+                        "rayQueryGenerateIntersection" => {
+                            let mut args = ctx.prepare_args(arguments, 2, span);
+                            let query = self.ray_query_pointer(args.next()?, ctx)?;
+                            let hit_t = self.expression(args.next()?, ctx)?;
+                            args.finish()?;
+
+                            let fun = crate::RayQueryFunction::GenerateIntersection { hit_t };
+                            let rctx = ctx.runtime_expression_ctx(span)?;
+                            rctx.block
+                                .push(crate::Statement::RayQuery { query, fun }, span);
+                            return Ok(None);
+                        }
+                        "rayQueryConfirmIntersection" => {
+                            let mut args = ctx.prepare_args(arguments, 1, span);
+                            let query = self.ray_query_pointer(args.next()?, ctx)?;
+                            args.finish()?;
+
+                            let fun = crate::RayQueryFunction::ConfirmIntersection;
+                            let rctx = ctx.runtime_expression_ctx(span)?;
+                            rctx.block
+                                .push(crate::Statement::RayQuery { query, fun }, span);
+                            return Ok(None);
+                        }
+                        "rayQueryTerminate" => {
+                            let mut args = ctx.prepare_args(arguments, 1, span);
+                            let query = self.ray_query_pointer(args.next()?, ctx)?;
+                            args.finish()?;
+
+                            let fun = crate::RayQueryFunction::Terminate;
+                            let rctx = ctx.runtime_expression_ctx(span)?;
+                            rctx.block
+                                .push(crate::Statement::RayQuery { query, fun }, span);
+                            return Ok(None);
+                        }
                         "rayQueryGetCommittedIntersection" => {
                             let mut args = ctx.prepare_args(arguments, 1, span);
                             let query = self.ray_query_pointer(args.next()?, ctx)?;
diff --git a/naga/src/lib.rs b/naga/src/lib.rs
index 04f3f52bba..2e917d34e0 100644
--- a/naga/src/lib.rs
+++ b/naga/src/lib.rs
@@ -1779,6 +1779,16 @@ pub enum RayQueryFunction {
         result: Handle<Expression>,
     },
 
+    /// Add a candidate generated intersection to be included
+    /// in the determination of the closest hit for a ray query.
+    GenerateIntersection {
+        hit_t: Handle<Expression>,
+    },
+
+    /// Confirm a triangle intersection to be included in the determination of
+    /// the closest hit for a ray query.
+    ConfirmIntersection,
+
     Terminate,
 }
 
diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs
index 8417bf77be..eb00a15790 100644
--- a/naga/src/valid/analyzer.rs
+++ b/naga/src/valid/analyzer.rs
@@ -1080,13 +1080,20 @@ impl FunctionInfo {
                 }
                 S::RayQuery { query, ref fun } => {
                     let _ = self.add_ref(query);
-                    if let crate::RayQueryFunction::Initialize {
-                        acceleration_structure,
-                        descriptor,
-                    } = *fun
-                    {
-                        let _ = self.add_ref(acceleration_structure);
-                        let _ = self.add_ref(descriptor);
+                    match *fun {
+                        crate::RayQueryFunction::Initialize {
+                            acceleration_structure,
+                            descriptor,
+                        } => {
+                            let _ = self.add_ref(acceleration_structure);
+                            let _ = self.add_ref(descriptor);
+                        }
+                        crate::RayQueryFunction::Proceed { result: _ } => {}
+                        crate::RayQueryFunction::GenerateIntersection { hit_t } => {
+                            let _ = self.add_ref(hit_t);
+                        }
+                        crate::RayQueryFunction::ConfirmIntersection => {}
+                        crate::RayQueryFunction::Terminate => {}
                     }
                     FunctionUniformity::new()
                 }
diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs
index 9f92e708ce..eb74ca237b 100644
--- a/naga/src/valid/function.rs
+++ b/naga/src/valid/function.rs
@@ -177,6 +177,8 @@ pub enum FunctionError {
     InvalidRayDescriptor(Handle<crate::Expression>),
     #[error("Ray Query {0:?} does not have a matching type")]
     InvalidRayQueryType(Handle<crate::Type>),
+    #[error("Hit distance {0:?} must be an f32")]
+    InvalidHitDistanceType(Handle<crate::Expression>),
     #[error("Shader requires capability {0:?}")]
     MissingCapability(super::Capabilities),
     #[error(
@@ -1487,6 +1489,19 @@ impl super::Validator {
                         crate::RayQueryFunction::Proceed { result } => {
                             self.emit_expression(result, context)?;
                         }
+                        crate::RayQueryFunction::GenerateIntersection { hit_t } => {
+                            match *context.resolve_type(hit_t, &self.valid_expression_set)? {
+                                Ti::Scalar(crate::Scalar {
+                                    kind: crate::ScalarKind::Float,
+                                    width: _,
+                                }) => {}
+                                _ => {
+                                    return Err(FunctionError::InvalidHitDistanceType(hit_t)
+                                        .with_span_static(span, "invalid hit_t"))
+                                }
+                            }
+                        }
+                        crate::RayQueryFunction::ConfirmIntersection => {}
                         crate::RayQueryFunction::Terminate => {}
                     }
                 }
diff --git a/naga/src/valid/handles.rs b/naga/src/valid/handles.rs
index 260d442c79..af6b7b47ea 100644
--- a/naga/src/valid/handles.rs
+++ b/naga/src/valid/handles.rs
@@ -707,6 +707,10 @@ impl super::Validator {
                     crate::RayQueryFunction::Proceed { result } => {
                         validate_expr(result)?;
                     }
+                    crate::RayQueryFunction::GenerateIntersection { hit_t } => {
+                        validate_expr(hit_t)?;
+                    }
+                    crate::RayQueryFunction::ConfirmIntersection => {}
                     crate::RayQueryFunction::Terminate => {}
                 }
                 Ok(())


Test patch:
diff --git a/naga/tests/in/ray-query.wgsl b/naga/tests/in/ray-query.wgsl
index 9f94356b83..5fcc2012a0 100644
--- a/naga/tests/in/ray-query.wgsl
+++ b/naga/tests/in/ray-query.wgsl
@@ -87,5 +87,11 @@ fn main_candidate() {
     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);
+    if (intersection.kind == RAY_QUERY_INTERSECTION_AABB) {
+        rayQueryGenerateIntersection(&rq, 10.0);
+    } else if (intersection.kind == RAY_QUERY_INTERSECTION_TRIANGLE) {
+        rayQueryConfirmIntersection(&rq);
+    } else {
+        rayQueryTerminate(&rq);
+    }
 }
diff --git a/naga/tests/out/hlsl/ray-query.hlsl b/naga/tests/out/hlsl/ray-query.hlsl
index 4815d27c9a..2f44332bda 100644
--- a/naga/tests/out/hlsl/ray-query.hlsl
+++ b/naga/tests/out/hlsl/ray-query.hlsl
@@ -150,6 +150,16 @@ void main_candidate()
     float3 dir_2 = float3(0.0, 1.0, 0.0);
     rq.TraceRayInline(acc_struct, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2).flags, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2).cull_mask, RayDescFromRayDesc_(ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2)));
     RayIntersection intersection_1 = GetCandidateIntersection(rq);
-    output.Store(0, asuint(uint((intersection_1.kind == 3u))));
-    return;
+    if ((intersection_1.kind == 3u)) {
+        rq.CommitProceduralPrimitiveHit(10.0);
+        return;
+    } else {
+        if ((intersection_1.kind == 1u)) {
+            rq.CommitNonOpaqueTriangleHit();
+            return;
+        } else {
+            rq.Abort();
+            return;
+        }
+    }
 }
diff --git a/naga/tests/out/msl/overrides-ray-query.msl b/naga/tests/out/msl/overrides-ray-query.msl
index d70011159b..50372c8283 100644
--- a/naga/tests/out/msl/overrides-ray-query.msl
+++ b/naga/tests/out/msl/overrides-ray-query.msl
@@ -38,7 +38,6 @@ kernel void main_(
         if (metal::all(loop_bound == uint2(4294967295u))) { break; }
         loop_bound += uint2(loop_bound.y == 4294967295u, 1u);
         bool _e31 = rq.ready;
-        rq.ready = false;
         if (_e31) {
         } else {
             break;
diff --git a/naga/tests/out/msl/ray-query.msl b/naga/tests/out/msl/ray-query.msl
index 58927b1f12..2456fec759 100644
--- a/naga/tests/out/msl/ray-query.msl
+++ b/naga/tests/out/msl/ray-query.msl
@@ -58,7 +58,6 @@ RayIntersection query_loop(
         if (metal::all(loop_bound == uint2(4294967295u))) { break; }
         loop_bound += uint2(loop_bound.y == 4294967295u, 1u);
         bool _e9 = rq_1.ready;
-        rq_1.ready = false;
         if (_e9) {
         } else {
             break;
@@ -93,7 +92,6 @@ kernel void main_(
 
 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);
@@ -105,6 +103,14 @@ kernel void main_candidate(
     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;
+    if (intersection_1.kind == 3u) {
+        return;
+    } else {
+        if (intersection_1.kind == 1u) {
+            return;
+        } else {
+            rq.ready = false;
+            return;
+        }
+    }
 }
diff --git a/naga/tests/out/spv/ray-query.spvasm b/naga/tests/out/spv/ray-query.spvasm
index d9a9edc984..32c2d0fae2 100644
--- a/naga/tests/out/spv/ray-query.spvasm
+++ b/naga/tests/out/spv/ray-query.spvasm
@@ -1,14 +1,14 @@
 ; SPIR-V
 ; Version: 1.4
 ; Generator: rspirv
-; Bound: 160
+; Bound: 166
 OpCapability Shader
 OpCapability RayQueryKHR
 OpExtension "SPV_KHR_ray_query"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
 OpEntryPoint GLCompute %123 "main" %15 %17
-OpEntryPoint GLCompute %143 "main_candidate" %15 %17
+OpEntryPoint GLCompute %143 "main_candidate" %15
 OpExecutionMode %123 LocalSize 1 1 1
 OpExecutionMode %143 LocalSize 1 1 1
 OpMemberDecorate %10 0 Offset 0
@@ -92,7 +92,8 @@ OpMemberDecorate %18 0 Offset 0
 %129 = OpConstantComposite  %4  %109 %107 %109
 %132 = OpTypePointer StorageBuffer %6
 %137 = OpTypePointer StorageBuffer %4
-%146 = OpConstantComposite  %12  %27 %28 %29 %30 %128 %129
+%145 = OpConstantComposite  %12  %27 %28 %29 %30 %128 %129
+%146 = OpConstant  %3  10.0
 %57 = OpFunction  %10  None %56
 %59 = OpFunctionParameter  %32
 %60 = OpLabel
@@ -226,21 +227,35 @@ OpFunctionEnd
 %142 = OpLabel
 %147 = OpVariable  %32  Function
 %144 = OpLoad  %5  %15
-%145 = OpAccessChain  %126  %17 %64
 OpBranch %148
 %148 = OpLabel
-%149 = OpCompositeExtract  %6  %146 0
-%150 = OpCompositeExtract  %6  %146 1
-%151 = OpCompositeExtract  %3  %146 2
-%152 = OpCompositeExtract  %3  %146 3
-%153 = OpCompositeExtract  %4  %146 4
-%154 = OpCompositeExtract  %4  %146 5
+%149 = OpCompositeExtract  %6  %145 0
+%150 = OpCompositeExtract  %6  %145 1
+%151 = OpCompositeExtract  %3  %145 2
+%152 = OpCompositeExtract  %3  %145 3
+%153 = OpCompositeExtract  %4  %145 4
+%154 = OpCompositeExtract  %4  %145 5
 OpRayQueryInitializeKHR %147 %144 %149 %150 %153 %151 %154 %152
 %155 = OpFunctionCall  %10  %57 %147
 %156 = OpCompositeExtract  %6  %155 0
 %157 = OpIEqual  %8  %156 %78
-%158 = OpSelect  %6  %157 %62 %64
-%159 = OpAccessChain  %132  %145 %64
-OpStore %159 %158
+OpSelectionMerge %158 None
+OpBranchConditional %157 %159 %160
+%159 = OpLabel
+OpRayQueryGenerateIntersectionKHR %147 %146
+OpReturn
+%160 = OpLabel
+%161 = OpCompositeExtract  %6  %155 0
+%162 = OpIEqual  %8  %161 %62
+OpSelectionMerge %163 None
+OpBranchConditional %162 %164 %165
+%164 = OpLabel
+OpRayQueryConfirmIntersectionKHR %147
+OpReturn
+%165 = OpLabel
+OpReturn
+%163 = OpLabel
+OpBranch %158
+%158 = OpLabel
 OpReturn
 OpFunctionEnd
\ No newline at end of file