From 3efa0f79c3a553826fb81b09d8b3819f097ee8db Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Sat, 1 Feb 2025 00:58:15 -0800 Subject: [PATCH] Implement ray query candidate intersection generation and confirmation --- CHANGELOG.md | 1 + etc/specs/ray_tracing.md | 36 +++++++++++++------ naga/src/back/dot/mod.rs | 7 ++++ naga/src/back/hlsl/writer.rs | 13 +++++++ naga/src/back/msl/writer.rs | 38 ++++++++++++++++---- naga/src/back/pipeline_constants.rs | 4 +++ naga/src/back/spv/instructions.rs | 13 +++++++ naga/src/back/spv/ray.rs | 13 +++++++ naga/src/compact/statements.rs | 8 +++++ naga/src/front/wgsl/lower/mod.rs | 34 ++++++++++++++++++ naga/src/lib.rs | 10 ++++++ naga/src/valid/analyzer.rs | 21 +++++++---- naga/src/valid/function.rs | 15 ++++++++ naga/src/valid/handles.rs | 4 +++ naga/tests/in/ray-query.wgsl | 8 ++++- naga/tests/out/hlsl/ray-query.hlsl | 14 ++++++-- naga/tests/out/msl/overrides-ray-query.msl | 1 - naga/tests/out/msl/ray-query.msl | 14 +++++--- naga/tests/out/spv/ray-query.spvasm | 41 +++++++++++++++------- 19 files changed, 250 insertions(+), 45 deletions(-) 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, 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) -> bool` +rayQueryProceed(rq: ptr) -> 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) -> RayIntersection @@ -71,19 +80,24 @@ rayQueryGetCommittedIntersection(rq: ptr) -> RayIntersectio rayQueryGetCandidateIntersection(rq: ptr) -> 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 Writer { 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>, 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, }, + /// Add a candidate generated intersection to be included + /// in the determination of the closest hit for a ray query. + GenerateIntersection { + hit_t: Handle, + }, + + /// 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), #[error("Ray Query {0:?} does not have a matching type")] InvalidRayQueryType(Handle), + #[error("Hit distance {0:?} must be an f32")] + InvalidHitDistanceType(Handle), #[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(()) 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(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