diff --git a/CHANGELOG.md b/CHANGELOG.md index ea6c7e7cda6..a4be03532ad 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -100,7 +100,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148] - Fix textureNumLevels in the GLSL backend. By @magcius in [#6483](https://github.com/gfx-rs/wgpu/pull/6483). - 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 diff --git a/examples/src/ray_cube_compute/shader.wgsl b/examples/src/ray_cube_compute/shader.wgsl index 43604a99704..79ee7ad7e5d 100644 --- a/examples/src/ray_cube_compute/shader.wgsl +++ b/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) { 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; diff --git a/examples/src/ray_scene/shader.wgsl b/examples/src/ray_scene/shader.wgsl index f6bd2398c8b..4e16bd94535 100644 --- a/examples/src/ray_scene/shader.wgsl +++ b/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, diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index ed91aa24898..39bb2f0b21b 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -2254,14 +2254,14 @@ impl Writer { 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}(")?; diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index def4054dc09..3dc87b28d42 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -1736,10 +1736,7 @@ impl<'w> BlockContext<'w> { } 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) } }; diff --git a/naga/src/back/spv/ray.rs b/naga/src/back/spv/ray.rs index c2daf4b3f64..7f16f803e41 100644 --- a/naga/src/back/spv/ray.rs +++ b/naga/src/back/spv/ray.rs @@ -106,23 +106,60 @@ impl<'w> BlockContext<'w> { &mut self, query: Handle, 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<'w> BlockContext<'w> { 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 { diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 111d4f18d99..030207b43d7 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/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( diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs index 50d10727576..d0288dabb13 100644 --- a/naga/src/front/wgsl/parse/mod.rs +++ b/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(); diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 0bbf857d0b5..dcab9d40e90 100644 --- a/naga/src/lib.rs +++ b/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 aren't expected to match. +#[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, } diff --git a/naga/tests/in/ray-query.wgsl b/naga/tests/in/ray-query.wgsl index ea38e65c17f..9f94356b830 100644 --- a/naga/tests/in/ray-query.wgsl +++ b/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(0.0); + let dir = vec3(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); +} diff --git a/naga/tests/out/msl/ray-query.msl b/naga/tests/out/msl/ray-query.msl index e5e44d0b5f7..b8230fb2e81 100644 --- a/naga/tests/out/msl/ray-query.msl +++ b/naga/tests/out/msl/ray-query.msl @@ -46,16 +46,16 @@ 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; @@ -63,7 +63,7 @@ RayIntersection query_loop( #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(intersection_1.kind == 3u); + return; +} diff --git a/naga/tests/out/spv/ray-query.spvasm b/naga/tests/out/spv/ray-query.spvasm index 8b784f2fa0c..5279bfc2e12 100644 --- a/naga/tests/out/spv/ray-query.spvasm +++ b/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 \ No newline at end of file